# This seems to be a false positive
src/gromacs/nbnxm/cuda/nbnxm_cuda_types.h: error: NbnxmGpu: is in internal file(s), but appears in public documentation
+# False positive caused by forward-declaration of BasicVector in src/testutils/refdata.h
+src/gromacs/math/vectypes.h: error: gmx::BasicVector: is in library file(s), but appears in public documentation
+
# Temporary while we change the SIMD implementation
src/gromacs/simd/impl_sparc64_hpc_ace/impl_sparc64_hpc_ace_common.h: warning: should include "simd.h"
.. figure:: plots/free2.*
:width: 6.00000cm
- Free energy cycles. **B:** to calculate
+ Free energy cycles. **B:** to calculate
:math:`\Delta G_{12}`, the free energy difference for binding of
inhibitors **I** respectively **I**\ :math:`^{\prime}` to enzyme
**E**.
*
* Output data contains a column for each data set in the input data, and a
* frame for each column in the input data. If different data sets have
- * different number of columns, the frame count accomodates the largest data
+ * different number of columns, the frame count accommodates the largest data
* set. Other columns are padded with zero values that are additionally marked
* as missing.
* Each value in the output data is the average of the corresponding
class DensityFittingForceProvider::Impl
{
public:
- //! \copydoc DensityFittingForceProvider(const DensityFittingParameters ¶meters)
+ //! \copydoc DensityFittingForceProvider::DensityFittingForceProvider
Impl(const DensityFittingParameters& parameters,
basic_mdspan<const float, dynamicExtents3D> referenceDensity,
const TranslateAndScale& transformationToDensityLattice,
else
{
const float3* __restrict__ gm_coordinates = asFloat3(kernelParams.atoms.d_coordinates);
- /* Recaclulate Splines */
+ /* Recalculate Splines */
if (c_useAtomDataPrefetch)
{
// charges
/*! \brief
* CUDA only
* Atom limit above which it is advantageous to turn on the
- * recalcuating of the splines in the gather and using less threads per atom in the spline and spread
+ * recalculating of the splines in the gather and using less threads per atom in the spline and spread
*/
constexpr int c_pmeGpuPerformanceAtomLimit = 23000;
* This is intended for use where there is no way to produce a data
* structure that would prevent execution from segfaulting. */
[[noreturn]] void warning_error_and_exit(warninp_t wi, const char* s, int f_errno, const char* file, int line);
-//! \copydoc warning_error_and_exit(warninp_t, const char *, int, const char *, int);
+//! \copydoc warning_error_and_exit(warninp_t,const char *,int,const char *,int)
[[noreturn]] void
warning_error_and_exit(warninp_t wi, const std::string& s, int f_errno, const char* file, int line);
*
* \tparam ValueType Raw value type of the \p buffer.
* \param[in,out] buffer Pointer to the device-side buffer.
- * \param[in] numValues Number of values to accomodate.
+ * \param[in] numValues Number of values to accommodate.
* \param[in] deviceContext The buffer's dummy device context - not managed explicitly in CUDA RT.
*/
template<typename ValueType>
*
* \tparam ValueType Raw value type of the \p buffer.
* \param[in,out] buffer Pointer to the device-side buffer.
- * \param[in] numValues Number of values to accomodate.
+ * \param[in] numValues Number of values to accommodate.
* \param[in] deviceContext The buffer's device context-to-be.
*/
template<typename ValueType>
}
}
- /* Caclulate GMX_SIMD_REAL_WIDTH dihedral angles at once */
+ /* Calculate GMX_SIMD_REAL_WIDTH dihedral angles at once */
dih_angle_simd(x, ai, aj, ak, al, pbc_simd, &phi_S, &mx_S, &my_S, &mz_S, &nx_S, &ny_S,
&nz_S, &nrkj_m2_S, &nrkj_n2_S, &p_S, &q_S);
}
}
- /* Caclulate GMX_SIMD_REAL_WIDTH dihedral angles at once */
+ /* Calculate GMX_SIMD_REAL_WIDTH dihedral angles at once */
dih_angle_simd(x, ai, aj, ak, al, pbc_simd, &phi_S, &mx_S, &my_S, &mz_S, &nx_S, &ny_S,
&nz_S, &nrkj_m2_S, &nrkj_n2_S, &p_S, &q_S);
* \param[in] kernelShapeParameters determine the shape of the spreading kernel
*/
explicit Impl(const GaussianSpreadKernelParameters::Shape& kernelShapeParameters);
- //! \copydoc DensityFittingForce::evaluateForce(const DensitySpreadKernelParameters::PositionAndAmplitude & localParameters, basic_mdspan<const float, dynamicExtents3D> densityDerivative)
+ //! \copydoc DensityFittingForce::evaluateForce
RVec evaluateForce(const GaussianSpreadKernelParameters::PositionAndAmplitude& localParameters,
basic_mdspan<const float, dynamicExtents3D> densityDerivative);
//! The width of the Gaussian in lattice spacing units
//! \}
PaddedVector() : storage_(), unpaddedEnd_(begin()) {}
- /*! \brief Constructor that specifes the initial size. */
+ /*! \brief Constructor that specifies the initial size. */
explicit PaddedVector(size_type count, const allocator_type& allocator = Allocator()) :
storage_(count, allocator),
unpaddedEnd_(begin() + count)
// the padding elements are added
resizeWithPadding(count);
}
- /*! \brief Constructor that specifes the initial size and an element to copy. */
+ /*! \brief Constructor that specifies the initial size and an element to copy. */
explicit PaddedVector(size_type count, value_type const& v, const allocator_type& allocator = Allocator()) :
storage_(count, v, allocator),
unpaddedEnd_(begin() + count)
{
return ::norm2(v);
}
+/*! \} */
} // namespace gmx
-/*! \} */
#endif
/*!\brief Number of CUDA threads in a block
*
- * \todo Check if using smaller block size will lead to better prformance.
+ * \todo Check if using smaller block size will lead to better performance.
*/
constexpr static int c_threadsPerBlock = 256;
//! Maximum number of threads in a block (for __launch_bounds__)
{
/*!\brief Number of CUDA threads in a block
*
- * \todo Check if using smaller block size will lead to better prformance.
+ * \todo Check if using smaller block size will lead to better performance.
*/
constexpr static int c_threadsPerBlock = 256;
//! Maximum number of threads in a block (for __launch_bounds__)
launchGpuKernel(scaleCoordinates_kernel, coordinateScalingKernelLaunchConfig_, deviceStream_,
nullptr, "scaleCoordinates_kernel", kernelArgs);
// TODO: Although this only happens on the pressure coupling steps, this synchronization
- // can affect the perfornamce if nstpcouple is small.
+ // can affect the performance if nstpcouple is small.
deviceStream_.synchronize();
}
launchGpuKernel(scaleCoordinates_kernel, coordinateScalingKernelLaunchConfig_, deviceStream_,
nullptr, "scaleCoordinates_kernel", kernelArgs);
// TODO: Although this only happens on the pressure coupling steps, this synchronization
- // can affect the perfornamce if nstpcouple is small.
+ // can affect the performance if nstpcouple is small.
deviceStream_.synchronize();
}
#include "pairlist.h"
-/*! \copybrief Returns the base-2 log of n.
+/*! \brief Returns the base-2 log of n.
* *
* Generates a fatal error when n is not an integer power of 2.
*/
* \param[in,out] gpu_nbv The nonbonded data GPU structure.
* \param[in] d_x Device-side coordinates in plain rvec format.
* \param[in] xReadyOnDevice Event synchronizer indicating that the coordinates are ready in
- * the device memory. \param[in] locality Copy coordinates for local or non-local atoms.
+ * the device memory.
+ * \param[in] locality Copy coordinates for local or non-local atoms.
* \param[in] gridId Index of the grid being converted.
* \param[in] numColumnsMax Maximum number of columns in the grid.
*/
tables.tableF.size(), deviceContext);
}
-void inline printEnviromnentVariableDeprecationMessage(bool isEnvironmentVariableSet,
+void inline printEnvironmentVariableDeprecationMessage(bool isEnvironmentVariableSet,
const std::string& environmentVariableSuffix)
{
if (isEnvironmentVariableSet)
const bool forceTwinCutoffEwaldLegacy = (getenv("GMX_CUDA_NB_EWALD_TWINCUT") != nullptr)
|| (getenv("GMX_OCL_NB_EWALD_TWINCUT") != nullptr);
- printEnviromnentVariableDeprecationMessage(forceAnalyticalEwaldLegacy, "NB_ANA_EWALD");
- printEnviromnentVariableDeprecationMessage(forceTabulatedEwaldLegacy, "NB_TAB_EWALD");
- printEnviromnentVariableDeprecationMessage(forceTwinCutoffEwaldLegacy, "NB_EWALD_TWINCUT");
+ printEnvironmentVariableDeprecationMessage(forceAnalyticalEwaldLegacy, "NB_ANA_EWALD");
+ printEnvironmentVariableDeprecationMessage(forceTabulatedEwaldLegacy, "NB_TAB_EWALD");
+ printEnvironmentVariableDeprecationMessage(forceTwinCutoffEwaldLegacy, "NB_EWALD_TWINCUT");
const bool forceAnalyticalEwald =
(getenv("GMX_GPU_NB_ANA_EWALD") != nullptr) || forceAnalyticalEwaldLegacy;
//! Define 1/sqrt(pi)
#define M_FLOAT_1_SQRTPI 0.564189583547756f
-/*! @} */
/*! \brief Constants for platform-dependent defaults for the prune kernel's j4 processing concurrency.
*
* Initialized using macros that can be overridden at compile-time (using #GMX_NBNXN_PRUNE_KERNEL_J4_CONCURRENCY).
/*!
* \brief Wrap a restraint potential as an MDModule
*
- * \param restraint shared ownership of an object for calculating restraint forces.
- * \return new wrapper object sharing ownership of restraint.
- *
* Consumers of the interfaces provided by an IMDModule do not extend the lifetime
* of the interface objects returned by mdpOptionProvider(), outputProvider(), or
* registered via initForceProviders(). Calling code must keep this object alive
* as long as those interfaces are needed (probably the duration of an MD run).
*
- * \param restraint handle to object to wrap
+ * \param restraint shared ownership of an object for calculating restraint forces
* \param sites list of sites for the framework to pass to the restraint
+ * \return new wrapper object sharing ownership of restraint
*/
static std::unique_ptr<RestraintMDModule> create(std::shared_ptr<gmx::IRestraintPotential> restraint,
const std::vector<int>& sites);
/*
* This file is part of the GROMACS molecular simulation package.
*
- * Copyright (c) 2014,2015,2016,2017,2019, by the GROMACS development team, led by
+ * Copyright (c) 2014,2015,2016,2017,2019,2020, by the GROMACS development team, led by
* Mark Abraham, David van der Spoel, Berk Hess, and Erik Lindahl,
* and including many others, as listed in the AUTHORS file in the
* top-level source directory and at http://www.gromacs.org.
/*! \addtogroup module_simd */
/*! \{ */
-/* \name SIMD implementation data types
+/*! \name SIMD implementation data types
* \{
*/
/*
* This file is part of the GROMACS molecular simulation package.
*
- * Copyright (c) 2014,2015,2016,2017,2019, by the GROMACS development team, led by
+ * Copyright (c) 2014,2015,2016,2017,2019,2020, by the GROMACS development team, led by
* Mark Abraham, David van der Spoel, Berk Hess, and Erik Lindahl,
* and including many others, as listed in the AUTHORS file in the
* top-level source directory and at http://www.gromacs.org.
/*! \addtogroup module_simd */
/*! \{ */
-/* \name SIMD implementation data types and built-in conversions between types
+/*! \name SIMD implementation data types and built-in conversions between types
* \{
*/
/*! \addtogroup module_simd */
/*! \{ */
-/* \name Higher-level SIMD utility functions, double precision.
+/*! \name Higher-level SIMD utility functions, double precision.
*
* These include generic functions to work with triplets of data, typically
* coordinates, and a few utility functions to load and update data in the
/*! \addtogroup module_simd */
/*! \{ */
-/* \name Higher-level SIMD utility functions, single precision.
+/*! \name Higher-level SIMD utility functions, single precision.
*
* These include generic functions to work with triplets of data, typically
* coordinates, and a few utility functions to load and update data in the
namespace internal
{
-// TODO: Don't foward function but properly rename them and use proper traits
+// TODO: Don't forward function but properly rename them and use proper traits
template<typename T>
struct Simd4Traits
{
} // namespace gmx
-// \} end of module_simd
+//! \} end of module_simd
//! \endcond end of condition libapi
/*! \brief SIMD pow(x,y), only targeting single accuracy.
*
- * \copydetails pow(SimdFloat)
+ * \copydetails pow(SimdFloat,SimdFloat)
*/
template<MathOptimization opt = MathOptimization::Safe>
static inline SimdFloat gmx_simdcall powSingleAccuracy(SimdFloat x, SimdFloat y)
// Actual math function tests below
-
-namespace
-{
-
/*! \cond internal */
/*! \addtogroup module_simd */
/*! \{ */
+namespace
+{
+
// Reference data is selected based on test name, so make the test name precision-dependent
# if GMX_DOUBLE
TEST_F(SimdMathTest, generateTestPointsDouble)
iterator end_;
};
-//! \copydoc ArrayRef::fromArray()
+/*! \brief
+ * Constructs a reference to a C array.
+ *
+ * \param[in] begin Pointer to the beginning of array.
+ * \param[in] size Number of elements in array.
+ *
+ * Passed array must remain valid for the lifetime of this object.
+ */
//! \related ArrayRef
template<typename T>
ArrayRef<T> arrayRefFromArray(T* begin, size_t size)
return ArrayRef<T>(begin, begin + size);
}
-//! \copydoc ArrayRef::fromArray()
+//! \copydoc arrayRefFromArray
//! \related ArrayRef
template<typename T>
ArrayRef<const T> constArrayRefFromArray(const T* begin, size_t size)