Artem Zhmurov [Fri, 2 Apr 2021 12:31:27 +0000 (15:31 +0300)]
Add early returns in clearDeviceBuffer(...) in OpenCL and CUDA
Clearing of empty buffer can cause the OpenCL/CUDA API failure.
Early return when requested size to clear is zero allows to
avoid this failure not calling the API for empty buffers.
This also makes the definition of the clearDeviceBuffer(...) in
OpenCL, CUDA similar to its definition in SYCL, where the early
return is already present.
Fixes #4002.
Joe Jordan [Sun, 11 Apr 2021 00:48:56 +0000 (00:48 +0000)]
Pass ref_t to initialize_lambdas
Dmitry Morozov [Wed, 7 Apr 2021 14:18:12 +0000 (17:18 +0300)]
Fix several errors when compiling with MSVC
ejjordan [Tue, 6 Apr 2021 09:21:22 +0000 (11:21 +0200)]
Minor forcerec cleanup
* gmx_bool -> bool
* rename a variable
* rvec -> RVec
* more array and vector
* now has default destructor
Kevin Boyd [Sat, 10 Apr 2021 20:25:12 +0000 (20:25 +0000)]
Clean up leaks
Artem Zhmurov [Sat, 3 Apr 2021 09:51:17 +0000 (12:51 +0300)]
Rename NBAtomData into NBAtomDataGpu
NBAtomData can only be present in GPU code, hence it should has
Gpu prefix in its name.
Paul Bauer [Fri, 2 Apr 2021 15:58:10 +0000 (15:58 +0000)]
Replace gmx_mtop_ilistloop_all
ejjordan [Thu, 1 Apr 2021 08:08:57 +0000 (10:08 +0200)]
Move domdec/partition into gmx namespace; minor cleanup
* Move functions in partition header into gmx namespace
* Localize variables
* Remove an in-out parameter that is never used
Andrey Alekseenko [Thu, 1 Apr 2021 13:40:30 +0000 (15:40 +0200)]
Ensure coordinates are copied for dipole moment calculation
With GPU update, we should wait for the coordinates to be copied to host
before using them for calculating the dipole moment.
This did not seem to cause any issues with CUDA but was causing a unit
test failure with SYCL (with !1329):
SYCL_BE=PI_OPENCL GMX_USE_GPU_BUFFER_OPS=1 GMX_FORCE_UPDATE_DEFAULT_GPU=1 ./bin/mdrun-test --gtest_filter=EwaldSurfaceTerm/EwaldSurfaceTermTest.WithinTolerances/0
Refs #3930, #3932
Paul Bauer [Thu, 1 Apr 2021 09:18:47 +0000 (11:18 +0200)]
Make FileOptionTypes enum class
Changed to enum class and made mapping to legacy types use
EnumerationArray.
Removed Unknown type as it wasn't used except for error checking. Used
Count type instead in its place.
Pascal Merz [Fri, 2 Apr 2021 00:53:32 +0000 (00:53 +0000)]
Clarify connection of temperature / pressure coupling and propagators
Temperature and pressure coupling and propagators are linked, as most
thermostats and barostats apply some type of scaling to positions and
/ or velocities. The temperature and pressure coupling objects must
therefore be linked to the propagators to properly function.
The current implementation did not allow for much flexibility, as it
only allowed propagators to decide whether to connect to a thermostat
or barostat or not. To allow for more complex algorithms, more
flexibility is required. To that end, this change introduces a
PropagatorTag object, which tags the propagators, and allows the
temperature and pressure coupling objects to pick which propagator to
act on. Additionally, this moves the responsibility of connecting to
the coupling algorithm rather than the propagator, which makes more
sense conceptually.
Refs #3423
Sebastian Kehl [Thu, 1 Apr 2021 14:38:08 +0000 (14:38 +0000)]
Add unit test for 1-4 interactions.
Paul Bauer [Thu, 1 Apr 2021 11:09:19 +0000 (13:09 +0200)]
Remove dead code in sasa module
If needed, the code can be found again in git.
Mark Abraham [Wed, 31 Mar 2021 08:31:48 +0000 (10:31 +0200)]
Make EnumerationArray::size() not static
It's formally correct that the method can be static, but this makes
code more complex at the point of use for the doubtful benefit of
making clear to the human reader that the value returned by size() is
a compile-time constant. The compiler can always see it.
Joe Jordan [Wed, 31 Mar 2021 16:10:35 +0000 (16:10 +0000)]
ArrayRefs and const refs in global_stat
Use const refs where possible in global_stat. Use ArrayRef for
signalBuffer and return an int instead of passing pointer to int.
Also movemost variable definitions to closer to their first use.
Berk Hess [Wed, 31 Mar 2021 13:54:14 +0000 (13:54 +0000)]
Use EnumerationArray for pull isAngleType
Alan Gray [Thu, 25 Feb 2021 13:23:11 +0000 (05:23 -0800)]
Remove PME coordinate send code duplication
Unifies codepath for PME coordinate send for default and GPU direct
codepaths. The only effect this has on ordering of activities is for a
run config with GPU buffer ops, without GPU direct communication and
without GPU update, where the (CPU-side) PME-PP coordinate transfer
will now be slightly delayed to allow for an H2D launch of coordinate
data.
Partly addresses #3914
Paul Bauer [Wed, 31 Mar 2021 11:18:01 +0000 (11:18 +0000)]
Move cluster methods to its own file
Breaks out some of the methods in gmx cluster into a separate file
before future refactoring.
Refs #4000
ejjordan [Wed, 31 Mar 2021 08:35:55 +0000 (10:35 +0200)]
Fix doxygen for updateMDLeapfrogGeneral
Fixes an issue created in !1350.
Joe Jordan [Wed, 31 Mar 2021 08:30:23 +0000 (08:30 +0000)]
Use more ArrayRefs in signatures of update impl class
Part of ongoing work to make refactoring mdatoms easier. This only
changes the impl class of update. A followup change can propagate
using ArrayRef to the update class itself, which will make it easier
to call update in both tests and in nblib.
Berk Hess [Tue, 30 Mar 2021 20:28:52 +0000 (22:28 +0200)]
Use using for ArrayRef and RVec in pull code
Artem Zhmurov [Tue, 30 Mar 2021 15:08:15 +0000 (18:08 +0300)]
Unify gpu_free(...) function in OpenCL, CUDA and SYCL versions of NBNXM
Refs #2608
Artem Zhmurov [Wed, 31 Mar 2021 00:30:33 +0000 (00:30 +0000)]
Fix conditional on when DtoH forces copy occur
d2d4a50b4c636c203028c5bff311924ec15e7825 introduced performance
regression with forces copied from device to host on each step.
This fixes the issue by reinstantiating proper condition on the
copy call.
Fixes #4001
Refs #2608
Berk Hess [Tue, 30 Mar 2021 20:21:13 +0000 (20:21 +0000)]
Move dynamic pull groups and pull nthreads
Moved dynamic pull groups into their respective coordinates.
Set the number of threads per pull group instead of globally.
Now all thread parallel regions run on 1 thread with less than 100
local atoms.
Artem Zhmurov [Tue, 30 Mar 2021 14:44:12 +0000 (17:44 +0300)]
Use default constructor for the device buffers that can be uninitialized
Some buffers in NBNXM are initialized conditionally, which posseses
a problem for the OpenCL kernel jit compilation, where these buffers
are used as arguments.
Refs #2608
Paul Bauer [Tue, 30 Mar 2021 12:43:24 +0000 (14:43 +0200)]
Fix some include guards in gmx ana
Clang tidy would complain about the defines starting with _.
Andrey Alekseenko [Tue, 30 Mar 2021 12:39:45 +0000 (12:39 +0000)]
Allow the use of only one backend at a time with SYCL-DPCPP
This is not a hard requirement; multiple backends can be used together
just fine.
But in DPCPP, the same physical device can appear as different virtual
devices provided by different backends (e.g., the same GPU can be accessible
via both OpenCL and L0).
Thus, using devices from two backends is more likely to be a user error than
the desired behavior. In this function, we choose the backend with the most compatible
devices. In case of a tie, we choose OpenCL (if present), or some arbitrary backend
among those with the most devices.
In hipSYCL, the problem of using the same device twice is unlikely to manifest.
It has (as of 2021-03-03) another issues: D2D copy between different backends
is not allowed. We don't use D2D in SYCL yet. Additionally, hipSYCL does not implement
the `sycl::platform::get_backend()` function.
Thus, we only do the backend filtering with DPCPP to keep the code
simple.
If needed, the `GMX_GPU_DISABLE_COMPATIBILITY_CHECK` might be used to
disable this limitation, or a backend-specific device filtering (e.g.,
`SYCL_BE` and `SYCL_DEVICE_FILTER`) can be used to only
provide the desired backend.
kanduri [Tue, 30 Mar 2021 11:53:51 +0000 (11:53 +0000)]
Turn t_forcerec.shift_vec into an std::vector of gmx::RVec
Change function signatures so that call sites have minimal changes
Berk Hess [Tue, 30 Mar 2021 10:15:43 +0000 (10:15 +0000)]
Add coordIndex to t_pull_coord
The pull coordinate index is an integral part of a pull coordinate
and can not be changed. Thus it can be part of t_pull_coord which
allows avoiding index passing in several functions.
Mark Abraham [Tue, 30 Mar 2021 09:44:29 +0000 (09:44 +0000)]
Move responsibility for checking bondeds in reverse topology
The reverse topology is a data structure distributed across all
domains that should refer exactly once to all bonded interactions
present in the global topology. We check that this is true at run
time, but need a complex implementation to avoid doing extra
collective communication that supports running what is effectively an
assertion that the reverse topology construction is correct.
Past attempts to improve this (e.g.
8b6c99a1) left the code in an
intermediate state where the responsibility for this check was mostly
implemented in do_md rather than compute_globals. That gave the
impression of compute_globals() returning an integer
parameter. However as above it's really an implementation detail of a
check of the reverse topology, and so the logic and memory for that
check should be in the domain decomposition module and object.
This change move the details of that check to the rest of the code
that handles the reverse topology. It renames several variables and
adds documentation for them.
This simplifies parts of the implementation of the modular simulator,
and prepares for implementing the proposal for #3887 by making all
outputs from global reduction local to the module that consumes them.
Refs #3421
Refs #3887
Joe Jordan [Mon, 29 Mar 2021 13:03:56 +0000 (13:03 +0000)]
Add tests for calcvir
Andrey Alekseenko [Mon, 29 Mar 2021 10:43:47 +0000 (13:43 +0300)]
Fix webpage build
MR !1235 (
dff1d7e309eb751b6b7585853501c544919db59f) had a missing empty
line, which upset Sphinx.
Alan Gray [Mon, 29 Mar 2021 11:12:25 +0000 (11:12 +0000)]
Avoid unnecessary CPU force buffer clearing
Only clear CPU force buffers on steps that actually use the buffers.
Partly addresses #3970
Joe Jordan [Mon, 29 Mar 2021 08:35:17 +0000 (08:35 +0000)]
Move nsgrid header/source from mdlib to domdec
The function get_nsgrid_boundaries is only called in domdec module
so it should live there. Removed an unused constant and moved another
to the source which is the only consumer. Also removed a todo which
no longer makes sense.
Kevin Boyd [Mon, 29 Mar 2021 07:47:18 +0000 (07:47 +0000)]
Remove unused LSAN suppressions
Some unknown refactoring has left these unneeded
Mark Abraham [Sat, 27 Mar 2021 20:36:35 +0000 (20:36 +0000)]
Update developer guide on error handling
Deferred to C++ core guidelines on most content, plus
some GROMACS-specific clarification.
Andrey Alekseenko [Sat, 27 Mar 2021 19:24:34 +0000 (19:24 +0000)]
Avoid enqueueing the same event multiple times in nbnxn_atomdata_x_to_nbat_x_gpu
With several non-local grids, we were calling
nbnxnInsertNonlocalGpuDependency(nb, NonLocal) in nbnxn_gpu_x_to_nbat_x
after enqueueing the transformation kernel for each one.
Now, we only call it once after submitting all kernels for
given locality (marking the event for Local, enqueueing the wait for it
for NonLocal).
Refs #3988.
Paul Bauer [Wed, 24 Mar 2021 10:17:11 +0000 (11:17 +0100)]
Change logic for checksumming availability
Changed order of checks to see if we can calculate the checksum or not.
Joe Jordan [Fri, 26 Mar 2021 12:08:14 +0000 (12:08 +0000)]
Use ArrayRefs instead of mdatoms in gmx_nb_free_energy_kernel signature
In addition to helping facilitate the coming mdatoms refactoring,
this is needed to facilitate calling gmx_nb_free_energy_kernel from
nblib without constructing mdatoms.
Szilárd Páll [Fri, 26 Mar 2021 08:14:16 +0000 (08:14 +0000)]
Allow setting nbnxm cluster size for SYCL
Rename GMX_OPENCL_NB_CLUSTER_SIZE to GMX_GPU_NB_CLUSTER_SIZE and use its
value in SYCL builds too. The former is still taken into account at
cmake-time, but users are advised to use the latter.
Refs #3847 #3933 #3935
Berk Hess [Thu, 25 Mar 2021 19:24:23 +0000 (19:24 +0000)]
Localize variables in do_lbfgs()
Also replaced pointers by std::vector.
Joe Jordan [Thu, 25 Mar 2021 19:04:20 +0000 (19:04 +0000)]
Use more vectors for EnergyOutput private members
Artem Zhmurov [Thu, 25 Mar 2021 18:28:42 +0000 (18:28 +0000)]
Unify gpu_upload_shiftvec(...) function in NBNXM
Refs. #2608
Joe Jordan [Thu, 25 Mar 2021 16:53:45 +0000 (16:53 +0000)]
Constructor for t_tabledata
Added default and specialized constructors for t_tabledata. With
one additional minor change this means that smalloc header is no
longer needed in forcetable.cpp.
ejjordan [Wed, 24 Mar 2021 22:33:40 +0000 (23:33 +0100)]
ArrayRef and const ref in gmx_nb_free_energy_kernel
Use const ref where possible in nonbonded free energy kernels. Also
use ArrayRef for passing coordinates.
Joe Jordan [Thu, 25 Mar 2021 10:48:38 +0000 (10:48 +0000)]
Use ArrayRef in special forces
Pass ArrayRef<RVec> instead of rvec* in edsam, pull_rotation, and
imd. Also removed unneeded header basedefinitions.
A followup change can refactor communicate_group_positions to take
ArrayRefs.
ejjordan [Wed, 24 Mar 2021 23:02:22 +0000 (00:02 +0100)]
Fix some doxygen errors
Mark Abraham [Thu, 25 Mar 2021 06:32:24 +0000 (07:32 +0100)]
Fixed unused variable warning
Joe Jordan [Thu, 25 Mar 2021 06:11:17 +0000 (06:11 +0000)]
Use ArrayRef in signatures of constraints and some pulling
In order to use ArrayRef in constr, settle, shake, and lincs, some
pull code also needed to be refactored. Part of preparation for
refactoring of t_mdatoms.
Paul Bauer [Wed, 24 Mar 2021 12:49:52 +0000 (12:49 +0000)]
Inline wallcycle counting functions
Inlines start and stop functions for wallcycle counting.
Andrey Alekseenko [Wed, 24 Mar 2021 11:54:55 +0000 (11:54 +0000)]
Unify nbnxn_gpu_init_x_to_nbat_x
Previously, we had this function only for CUDA, but we will need it for
SYCL (#3932). OpenCL implementation is unlikely to be needed, but should
not hurt either.
Refs. #2608
Andrey Alekseenko [Wed, 24 Mar 2021 10:52:51 +0000 (13:52 +0300)]
Add missing GMX_ASSERT(haveCopiedXFromGpu)
We make this assertion every time we call
stateGpu->waitCoordinatesReadyOnHost(AtomLocality::Local) in the
do_force function.
This does not fix any bugs, just makes the code a bit more consistent in
how it checks self-consistency.
Refs #3988.
ejjordan [Wed, 24 Mar 2021 08:14:07 +0000 (09:14 +0100)]
Remove nb_kernel_data_t
All data was copied from gmx_enerdata_t. Now use ArrayRefs to pass
the parameters to the kernels.
ejjordan [Tue, 23 Mar 2021 15:56:26 +0000 (16:56 +0100)]
Remove unneeded members from nb_kernel_data_t
Several unused members are removed from nb_kernel_data_t. Additionally,
some members that can be passed directly from the gmx_enerdata_t
struct are removed, and instead added to the call signature of
gmx_nb_free_energy_kernel. This expansion of the call signature
is the only way to get rid of the real* members of nb_kernel_data_t.
Also, function with no implementation is removed.
Paul Bauer [Tue, 23 Mar 2021 15:07:16 +0000 (16:07 +0100)]
Use TextWriter in msd tests
Removes raw file pointer for mdp input file writing.
Fixes #3986
Berk Hess [Tue, 23 Mar 2021 14:47:26 +0000 (14:47 +0000)]
Add comparison operator to BasicVector
The default comparison seemed to compare the memory addresses
and not the elements.
Szilárd Páll [Tue, 23 Mar 2021 13:44:45 +0000 (13:44 +0000)]
Fix deviceIdsAssigned() to return a unique list of IDs
The method originally returned the rank's task-to-device mapping
rather than the unique list of devices that have been assigned tasks.
When #ranks > #devices this list contains duplicate device IDs.
Due to the changes made in
e2a2fe80 the vector returned by
deviceIdsAssigned() was passed to setupGpuDevicePeerAccess()
which results in redundant attempts to repeatedly enable peer access for
the same device pairs.
Refs #3980
ejjordan [Tue, 23 Mar 2021 10:29:00 +0000 (11:29 +0100)]
Forward declare ArrayRef more and inlcude basedefinitions where needed
Remove many includes of ArrayRef that can be forward declared. Also
remove includes of basedefinitions in many headers. In some cases
removing basedefinitinos meant also changing gmx_bool to bool.
Note that no change to source files is needed, apart from possibly
adding the basedefintions include, because gmx_bool uses true and
false under the hood.
Joe Jordan [Tue, 23 Mar 2021 10:15:33 +0000 (10:15 +0000)]
Modernize t_nrnb
* Use array instead of pointer.
* Remove unused functions.
* Add doxygen to inc_nrnb.
* Use inline function instead of macro in all places where possible.
* Remove unused headers.
Joe Jordan [Tue, 23 Mar 2021 09:40:41 +0000 (09:40 +0000)]
Use ArrayRef in calc_mu
ejjordan [Mon, 22 Mar 2021 09:30:50 +0000 (10:30 +0100)]
More const ref in shellfc
Pass a const ref to t_mdatoms in shellfc functions. Part of work
to make refactoring t_mdatoms easier.
Gaurav Garg [Wed, 17 Mar 2021 08:27:35 +0000 (08:27 +0000)]
Detect CUDA-aware support in underlying MPI implementation
Checks CUDA-aware MPI support at cmake configuration time and at runtime. Print CUDA-aware MPI support status with gmx -version.
ejjordan [Mon, 22 Mar 2021 11:49:54 +0000 (12:49 +0100)]
Use ArrayRef in do_ewald
Also removed a function with no implementation.
Mark Abraham [Mon, 22 Mar 2021 09:49:30 +0000 (09:49 +0000)]
Use StringToEnumValueConverter and enumValueToString more
Follow up from MR 1213
Andrey Alekseenko [Mon, 22 Mar 2021 09:09:33 +0000 (09:09 +0000)]
Fix compiler and linter warnings
- Unused variables in pme_pp.cpp
- Insufficient NOLINTing in update.cpp
Andrey Alekseenko [Fri, 19 Mar 2021 21:55:19 +0000 (00:55 +0300)]
Remove incorrect assertion for GPU timing in SYCL
- The expression was backwards: it was failing when timing was disabled
instead of failing when it is enabled.
- The expression was run on a freshly-initialized data structure, and
the value of bDoTime was always overridden later, so it was not
actually checking anything.
Since we anyway never set bDoTime to true in SYCL build, we might as
well remove this assertion here altogether.
Introduced in !1296 (
c4a672b93b655ff674b671fd10ebe728b2d19ec8).
Kevin Boyd [Sat, 20 Mar 2021 21:43:58 +0000 (14:43 -0700)]
Clean up leaks in t_inputrec
Refs #3984
Szilárd Páll [Fri, 19 Mar 2021 19:04:42 +0000 (20:04 +0100)]
Remove a misplaced update cycle closing
This was removed in
b80b7a327b1 but got added back later presumably
during a merge.
Refs #3764
Szilárd Páll [Wed, 10 Mar 2021 10:50:47 +0000 (11:50 +0100)]
Rename sendFToPpCudaDirect
sendFSynchronizerToPpCudaDirect expresses what the method actually does.
Artem Zhmurov [Thu, 18 Mar 2021 06:49:18 +0000 (09:49 +0300)]
Move the atomToInteractionLocality(...) into locality.h
The atomToInteractionLocality(...) does not depend on NBNXM
and can naturally reside alongside the definitions of localitiy
enumerations. This allows to combine haveGpuShortRangeWork(...)
functions in NBNXM, avoding confuciton of having two very similar
functions with different signatures.
Andrey Alekseenko [Wed, 10 Mar 2021 11:53:49 +0000 (14:53 +0300)]
Store buffer without offset in GpuForceReduction::Impl::baseForce_
- Store the pointer to the beginning of the array and apply offset to
it when invoking the kernel. Offsets to other arrays are applied at
the invocation time, so it makes sense to do the same with baseForce_.
- This also avoid storing pointers to the middle of the device buffer,
which is not a good fit for OpenCL and SYCL.
This change is preparation for adding a SYCL implementation of GPU
update (Issue #3932).
Artem Zhmurov [Fri, 19 Mar 2021 16:48:38 +0000 (16:48 +0000)]
Unify init_gpu function in NBNXM
Refs. #2608
Andrey Alekseenko [Tue, 16 Mar 2021 10:03:24 +0000 (11:03 +0100)]
Use DeviceBuffer<RVec> in GPU force reduction and PME code
... instead of raw device pointers.
Preparation for #3932. PME change is incidental, the main focus is
GpuForceReduction.
Andrey Alekseenko [Tue, 16 Mar 2021 11:21:51 +0000 (12:21 +0100)]
Enable compiling LeapFrog with hipSYCL
No testing yet. Trivial change after !1268.
Closes #3941.
Paul Bauer [Fri, 19 Mar 2021 11:45:06 +0000 (12:45 +0100)]
Bulk change to const ref for mtop
Change function signatures in mtop_lookup.h and mtop_util.h to be const
ref only for mtop. Also change signatures in other files as I went
along.
This change is only refactoring.
Paul Bauer [Fri, 19 Mar 2021 07:59:47 +0000 (08:59 +0100)]
Make wallcycle comments proper doxygen
Should reduce noise in other refactoring commits.
Berk Hess [Fri, 19 Mar 2021 12:42:49 +0000 (12:42 +0000)]
Remove const cast in update code
Joe Jordan [Fri, 19 Mar 2021 11:45:23 +0000 (11:45 +0000)]
Use more arrayref in listed forces sigatures
Pass members of t_mdatoms directly in many places in listed_forces.
Part of work on making refactor of t_mdatoms easier.
Paul Bauer [Fri, 19 Mar 2021 11:11:49 +0000 (11:11 +0000)]
Remove mdatoms from forceprovider call signature
This makes it possible to test the forceprovider without having to build
a partial mdatoms datastructure first. Needed for future refactoring of
mdatoms.
Joe Jordan [Fri, 19 Mar 2021 10:50:48 +0000 (10:50 +0000)]
Use ArrayRef in ewald_LRcorrection
This will make refactoring t_mdatoms easier. Also changed
signatures of calculateLongRangeNonbondeds to use more
ArrayRef.
Joe Jordan [Fri, 19 Mar 2021 09:54:58 +0000 (09:54 +0000)]
Make non bonded energy terms enum class
Also renamed some of the enum values.
Andrey Alekseenko [Fri, 19 Mar 2021 08:13:39 +0000 (11:13 +0300)]
Fix a warning from CUDA 10.1
Discovered during nightly build:
warning: static_assert with no message is a C++17 extension [-Wc++17-extensions]
ejjordan [Thu, 18 Mar 2021 15:42:02 +0000 (16:42 +0100)]
Use ArrayRef in gmx_pme_send_parameters
This will make refactoring t_mdatoms easier.
ejjordan [Thu, 18 Mar 2021 16:05:02 +0000 (17:05 +0100)]
Use const refs in atoms2md signature
Const refs are now used for mtop and inputrec in atoms2md. Doxygen
is also added and inputrec renamed from ir.
Paul Bauer [Thu, 18 Mar 2021 17:12:34 +0000 (17:12 +0000)]
Add wallcycle counting tests
Basic tests for running main and sub counters, and to check overhead.
Andrey Alekseenko [Thu, 11 Mar 2021 10:48:06 +0000 (13:48 +0300)]
Use DeviceBuffer in GPU update and NBNXM code
... instead of raw device pointers and DeviceBuffer<float3>.
We try to use DeviceBuffer<Float3>, but in some places we have to use
DeviceVector<gmx::RVec>, until we can define FloatN types without
including any backend-specific headers. Currently, Float3 is defined as
gmx::RVec, so this should not cause any issues.
Also added some helper functions to convert RVec ̌<-> Float3 <-> float3.
Preparation for #3932 and #3941.
ejjordan [Wed, 17 Mar 2021 09:31:48 +0000 (10:31 +0100)]
Use ArrayRef<const real> in gmx_pme_do
Also changed some static functions to use ArrayRef. Some local
variables need to be vectors because they use operator[]. This
is part of preliminary work to make refactoring t_mdatoms
easier.
Paul Bauer [Wed, 17 Mar 2021 05:59:56 +0000 (05:59 +0000)]
Fix errors overlooked in gmx msd changes
Made some NOLINT sections more broad.
Fixed documentation.
Joe Jordan [Tue, 16 Mar 2021 18:03:19 +0000 (18:03 +0000)]
Use RVec and std::array in gmx_domdec_comm_t
Where possible, rvecs have been changed to RVecs and pointers
changed to std::arrays.
Kevin Boyd [Tue, 16 Mar 2021 17:28:24 +0000 (17:28 +0000)]
Migrate gmx msd to the trajectoryanalysis framework
Features to still be added
-tensor, for the MSD tensor
-pdb, for per molecule B-factors
-rmcomm, for system COM removal
-mw, for tuning mass-weighting options
-maxtau, for limiting resource usage and unneccessary computation (refs #3870)
-improvements to documentation and maybe some flag behavior changes (refs #3869)
Refs #2368
Artem Zhmurov [Tue, 16 Mar 2021 12:53:21 +0000 (15:53 +0300)]
Unify gpu_init_atomdata(...) function
Refs #2608
Artem Zhmurov [Tue, 16 Mar 2021 13:26:37 +0000 (16:26 +0300)]
Unify gpu_launch_cpyback(...) function in NBNXM
Refs #2608
Joe Jordan [Tue, 16 Mar 2021 11:21:10 +0000 (11:21 +0000)]
Remove conditional assignemnt of cell sizes from domdec partition
The variable cellsizesWithDlb is always initiallized since it is a
std::vector, so there is no reason to set it to nullptr in
domdec/partition.
Artem Zhmurov [Sat, 13 Mar 2021 13:13:14 +0000 (16:13 +0300)]
Unify gpu_init_atomdata(...) function
Refs #2608
Artem Zhmurov [Mon, 15 Mar 2021 16:57:02 +0000 (19:57 +0300)]
Move pmalloc(..)/pfree(...) to separate source files in CUDA/OpenCL/SYCL
Having pmalloc(..) and pfree(..) functions in unrelated files in CUDA,
OpenCL and SYCL makes it hard to use these functions in platform-agnostic
code. This creates common header and moves these functions to separate
source files.
Paul Bauer [Mon, 15 Mar 2021 13:56:18 +0000 (14:56 +0100)]
Add subcycle counting to MPI build
This code path has not been tested in CI and should be added.
Paul Bauer [Mon, 15 Mar 2021 14:59:03 +0000 (14:59 +0000)]
Move M_PI definition to math/units.h
This makes more sense to have the definition there with the other
constants.
Resolves build issues with MSVC not finding M_PI definition.
Needed some minor header fixes.
Paul Bauer [Mon, 15 Mar 2021 10:06:16 +0000 (11:06 +0100)]
Add another cool quote.
Recently seen on Twitter ...
Artem Zhmurov [Mon, 15 Mar 2021 09:12:04 +0000 (12:12 +0300)]
Do not overallocate shiftVec and fshifts buffers in OpenCL
When changing buffers from flat float format to Float3,
the sizes of the shiftVec and fshift buffers were not updated
in OpenCL NBNXM data management, which lead to the buffers being
overallocated. This fixes the issue.
Introduced in
c66827166fc9099ecd1a4a2f7080558df70bf529
Artem Zhmurov [Mon, 15 Mar 2021 10:16:45 +0000 (10:16 +0000)]
Add a cool quote from Elvis Presley's song