Rename UpdateConstraintsCuda to UpdateConstraintsGpu

This is to folow general naming conventions across the code.

Refs #3114.

Change-Id: Iec3678e1d4152d6442818fcd81020e021927757c

diff --git a/src/gromacs/mdlib/CMakeLists.txt b/src/gromacs/mdlib/CMakeLists.txt
index daffc92d96ce15359f9ccbabdd97fb9e4ab270ef..770028ba3b14d0e315ce66156b0c5672044f1f8a 100644 (file)
--- a/src/gromacs/mdlib/CMakeLists.txt
+++ b/src/gromacs/mdlib/CMakeLists.txt
@@ -44,7 +44,7 @@ if(GMX_USE_CUDA)
        leapfrog_gpu.cu
        lincs_gpu.cu
        settle_gpu.cu
-       update_constrain_cuda_impl.cu
+       update_constrain_gpu_impl.cu
        )
 endif()
 

diff --git a/src/gromacs/mdlib/update_constrain_cuda.h b/src/gromacs/mdlib/update_constrain_gpu.h
similarity index 92%
rename from src/gromacs/mdlib/update_constrain_cuda.h
rename to src/gromacs/mdlib/update_constrain_gpu.h
index 33e517bb1e9addf91c44bc68bef86aa8cfc28457..e0190c367f586d75680820dd3315871461405443 100644 (file)
--- a/src/gromacs/mdlib/update_constrain_cuda.h
+++ b/src/gromacs/mdlib/update_constrain_gpu.h
@@ -34,17 +34,15 @@
  */
 /*! \libinternal \file
  *
- * \brief Declaration of high-level functions of CUDA implementation of update and constrain class.
- *
- * \todo Change "cuda" suffix to "gpu"
+ * \brief Declaration of high-level functions of GPU implementation of update and constrain class.
  *
  * \author Artem Zhmurov <zhmurov@gmail.com>
  *
  * \ingroup module_mdlib
  * \inlibraryapi
  */
-#ifndef GMX_MDLIB_UPDATE_CONSTRAIN_CUDA_H
-#define GMX_MDLIB_UPDATE_CONSTRAIN_CUDA_H
+#ifndef GMX_MDLIB_UPDATE_CONSTRAIN_GPU_H
+#define GMX_MDLIB_UPDATE_CONSTRAIN_GPU_H
 
 #include "gromacs/gpu_utils/devicebuffer_datatype.h"
 #include "gromacs/mdtypes/group.h"
@@ -63,7 +61,7 @@ struct t_pbc;
 namespace gmx
 {
 
-class UpdateConstrainCuda
+class UpdateConstrainGpu
 {
 
 public:
@@ -82,12 +80,12 @@ public:
      * \param[in] commandStream     GPU stream to use. Can be nullptr.
      * \param[in] xUpdatedOnDevice  The event synchronizer to use to mark that update is done on the GPU.
      */
-    UpdateConstrainCuda(const t_inputrec&     ir,
-                        const gmx_mtop_t&     mtop,
-                        const void*           commandStream,
-                        GpuEventSynchronizer* xUpdatedOnDevice);
+    UpdateConstrainGpu(const t_inputrec&     ir,
+                       const gmx_mtop_t&     mtop,
+                       const void*           commandStream,
+                       GpuEventSynchronizer* xUpdatedOnDevice);
 
-    ~UpdateConstrainCuda();
+    ~UpdateConstrainGpu();
 
     /*! \brief Integrate
      *
@@ -158,7 +156,7 @@ public:
 
     /*! \brief
      * Returns whether the maximum number of coupled constraints is supported
-     * by the CUDA LINCS code.
+     * by the GPU LINCS code.
      *
      * \param[in] mtop The molecular topology
      */
@@ -171,4 +169,4 @@ private:
 
 } // namespace gmx
 
-#endif
+#endif // GMX_MDLIB_UPDATE_CONSTRAIN_GPU_H

diff --git a/src/gromacs/mdlib/update_constrain_cuda_impl.cpp b/src/gromacs/mdlib/update_constrain_gpu_impl.cpp
similarity index 60%
rename from src/gromacs/mdlib/update_constrain_cuda_impl.cpp
rename to src/gromacs/mdlib/update_constrain_gpu_impl.cpp
index 1fa2718d675ca5f343020f5fbb42c734713bb2ed..919ba9b36b97a81b5f86e0a8200752be1956321f 100644 (file)
--- a/src/gromacs/mdlib/update_constrain_cuda_impl.cpp
+++ b/src/gromacs/mdlib/update_constrain_gpu_impl.cpp
@@ -44,75 +44,75 @@
 
 #include "config.h"
 
-#include "gromacs/mdlib/update_constrain_cuda.h"
+#include "gromacs/mdlib/update_constrain_gpu.h"
 
 #if GMX_GPU != GMX_GPU_CUDA
 
 namespace gmx
 {
 
-class UpdateConstrainCuda::Impl
+class UpdateConstrainGpu::Impl
 {
 };
 
-UpdateConstrainCuda::UpdateConstrainCuda(const t_inputrec& /* ir   */,
-                                         const gmx_mtop_t& /* mtop */,
-                                         const void* /* commandStream */,
-                                         GpuEventSynchronizer* /* xUpdatedOnDevice */) :
+UpdateConstrainGpu::UpdateConstrainGpu(const t_inputrec& /* ir   */,
+                                       const gmx_mtop_t& /* mtop */,
+                                       const void* /* commandStream */,
+                                       GpuEventSynchronizer* /* xUpdatedOnDevice */) :
     impl_(nullptr)
 {
     GMX_ASSERT(false,
                "A CPU stub for UpdateConstrain was called instead of the correct implementation.");
 }
 
-UpdateConstrainCuda::~UpdateConstrainCuda() = default;
+UpdateConstrainGpu::~UpdateConstrainGpu() = default;
 
-void UpdateConstrainCuda::integrate(GpuEventSynchronizer* /* fReadyOnDevice */,
-                                    const real /* dt */,
-                                    const bool /* updateVelocities */,
-                                    const bool /* computeVirial */,
-                                    tensor /* virialScaled */,
-                                    const bool /* doTemperatureScaling */,
-                                    gmx::ArrayRef<const t_grp_tcstat> /* tcstat */,
-                                    const bool /* doParrinelloRahman */,
-                                    const float /* dtPressureCouple */,
-                                    const matrix /* prVelocityScalingMatrix*/)
+void UpdateConstrainGpu::integrate(GpuEventSynchronizer* /* fReadyOnDevice */,
+                                   const real /* dt */,
+                                   const bool /* updateVelocities */,
+                                   const bool /* computeVirial */,
+                                   tensor /* virialScaled */,
+                                   const bool /* doTemperatureScaling */,
+                                   gmx::ArrayRef<const t_grp_tcstat> /* tcstat */,
+                                   const bool /* doParrinelloRahman */,
+                                   const float /* dtPressureCouple */,
+                                   const matrix /* prVelocityScalingMatrix*/)
 {
     GMX_ASSERT(false,
                "A CPU stub for UpdateConstrain was called instead of the correct implementation.");
 }
 
-void UpdateConstrainCuda::scaleCoordinates(const matrix /* scalingMatrix */)
+void UpdateConstrainGpu::scaleCoordinates(const matrix /* scalingMatrix */)
 {
     GMX_ASSERT(false,
                "A CPU stub for UpdateConstrain was called instead of the correct implementation.");
 }
 
-void UpdateConstrainCuda::set(DeviceBuffer<float> /* d_x */,
-                              DeviceBuffer<float> /* d_v */,
-                              const DeviceBuffer<float> /* d_f */,
-                              const t_idef& /* idef */,
-                              const t_mdatoms& /* md */,
-                              const int /* numTempScaleValues */)
+void UpdateConstrainGpu::set(DeviceBuffer<float> /* d_x */,
+                             DeviceBuffer<float> /* d_v */,
+                             const DeviceBuffer<float> /* d_f */,
+                             const t_idef& /* idef */,
+                             const t_mdatoms& /* md */,
+                             const int /* numTempScaleValues */)
 {
     GMX_ASSERT(false,
                "A CPU stub for UpdateConstrain was called instead of the correct implementation.");
 }
 
-void UpdateConstrainCuda::setPbc(const PbcType /* pbcType */, const matrix /* box */)
+void UpdateConstrainGpu::setPbc(const PbcType /* pbcType */, const matrix /* box */)
 {
     GMX_ASSERT(false,
                "A CPU stub for UpdateConstrain was called instead of the correct implementation.");
 }
 
-GpuEventSynchronizer* UpdateConstrainCuda::getCoordinatesReadySync()
+GpuEventSynchronizer* UpdateConstrainGpu::getCoordinatesReadySync()
 {
     GMX_ASSERT(false,
                "A CPU stub for UpdateConstrain was called instead of the correct implementation.");
     return nullptr;
 }
 
-bool UpdateConstrainCuda::isNumCoupledConstraintsSupported(const gmx_mtop_t& /* mtop */)
+bool UpdateConstrainGpu::isNumCoupledConstraintsSupported(const gmx_mtop_t& /* mtop */)
 {
     return false;
 }

diff --git a/src/gromacs/mdlib/update_constrain_cuda_impl.cu b/src/gromacs/mdlib/update_constrain_gpu_impl.cu
similarity index 68%
rename from src/gromacs/mdlib/update_constrain_cuda_impl.cu
rename to src/gromacs/mdlib/update_constrain_gpu_impl.cu
index 067920cb41e827179b777c910b455a5a0e747043..ebbe15a02ff5059be0305a748f05651ad9fe4a33 100644 (file)
--- a/src/gromacs/mdlib/update_constrain_cuda_impl.cu
+++ b/src/gromacs/mdlib/update_constrain_gpu_impl.cu
@@ -47,7 +47,7 @@
  */
 #include "gmxpre.h"
 
-#include "update_constrain_cuda_impl.h"
+#include "update_constrain_gpu_impl.h"
 
 #include <assert.h>
 #include <stdio.h>
@@ -63,7 +63,7 @@
 #include "gromacs/mdlib/leapfrog_gpu.cuh"
 #include "gromacs/mdlib/lincs_gpu.cuh"
 #include "gromacs/mdlib/settle_gpu.cuh"
-#include "gromacs/mdlib/update_constrain_cuda.h"
+#include "gromacs/mdlib/update_constrain_gpu.h"
 
 namespace gmx
 {
@@ -102,16 +102,16 @@ __launch_bounds__(c_maxThreadsPerBlock) __global__
     }
 }
 
-void UpdateConstrainCuda::Impl::integrate(GpuEventSynchronizer*             fReadyOnDevice,
-                                          const real                        dt,
-                                          const bool                        updateVelocities,
-                                          const bool                        computeVirial,
-                                          tensor                            virial,
-                                          const bool                        doTemperatureScaling,
-                                          gmx::ArrayRef<const t_grp_tcstat> tcstat,
-                                          const bool                        doParrinelloRahman,
-                                          const float                       dtPressureCouple,
-                                          const matrix                      prVelocityScalingMatrix)
+void UpdateConstrainGpu::Impl::integrate(GpuEventSynchronizer*             fReadyOnDevice,
+                                         const real                        dt,
+                                         const bool                        updateVelocities,
+                                         const bool                        computeVirial,
+                                         tensor                            virial,
+                                         const bool                        doTemperatureScaling,
+                                         gmx::ArrayRef<const t_grp_tcstat> tcstat,
+                                         const bool                        doParrinelloRahman,
+                                         const float                       dtPressureCouple,
+                                         const matrix                      prVelocityScalingMatrix)
 {
     // Clearing virial matrix
     // TODO There is no point in having separate virial matrix for constraints
@@ -145,7 +145,7 @@ void UpdateConstrainCuda::Impl::integrate(GpuEventSynchronizer*             fRea
     return;
 }
 
-void UpdateConstrainCuda::Impl::scaleCoordinates(const matrix scalingMatrix)
+void UpdateConstrainGpu::Impl::scaleCoordinates(const matrix scalingMatrix)
 {
     ScalingMatrix mu;
     mu.xx = scalingMatrix[XX][XX];
@@ -164,10 +164,10 @@ void UpdateConstrainCuda::Impl::scaleCoordinates(const matrix scalingMatrix)
     gpuStreamSynchronize(commandStream_);
 }
 
-UpdateConstrainCuda::Impl::Impl(const t_inputrec&     ir,
-                                const gmx_mtop_t&     mtop,
-                                const void*           commandStream,
-                                GpuEventSynchronizer* xUpdatedOnDevice) :
+UpdateConstrainGpu::Impl::Impl(const t_inputrec&     ir,
+                               const gmx_mtop_t&     mtop,
+                               const void*           commandStream,
+                               GpuEventSynchronizer* xUpdatedOnDevice) :
     coordinatesReady_(xUpdatedOnDevice)
 {
     GMX_ASSERT(xUpdatedOnDevice != nullptr, "The event synchronizer can not be nullptr.");
@@ -186,14 +186,14 @@ UpdateConstrainCuda::Impl::Impl(const t_inputrec&     ir,
     coordinateScalingKernelLaunchConfig_.stream           = commandStream_;
 }
 
-UpdateConstrainCuda::Impl::~Impl() {}
+UpdateConstrainGpu::Impl::~Impl() {}
 
-void UpdateConstrainCuda::Impl::set(DeviceBuffer<float>       d_x,
-                                    DeviceBuffer<float>       d_v,
-                                    const DeviceBuffer<float> d_f,
-                                    const t_idef&             idef,
-                                    const t_mdatoms&          md,
-                                    const int                 numTempScaleValues)
+void UpdateConstrainGpu::Impl::set(DeviceBuffer<float>       d_x,
+                                   DeviceBuffer<float>       d_v,
+                                   const DeviceBuffer<float> d_f,
+                                   const t_idef&             idef,
+                                   const t_mdatoms&          md,
+                                   const int                 numTempScaleValues)
 {
     GMX_ASSERT(d_x != nullptr, "Coordinates device buffer should not be null.");
     GMX_ASSERT(d_v != nullptr, "Velocities device buffer should not be null.");
@@ -219,67 +219,67 @@ void UpdateConstrainCuda::Impl::set(DeviceBuffer<float>       d_x,
             (numAtoms_ + c_threadsPerBlock - 1) / c_threadsPerBlock;
 }
 
-void UpdateConstrainCuda::Impl::setPbc(const PbcType pbcType, const matrix box)
+void UpdateConstrainGpu::Impl::setPbc(const PbcType pbcType, const matrix box)
 {
     setPbcAiuc(numPbcDimensions(pbcType), box, &pbcAiuc_);
 }
 
-GpuEventSynchronizer* UpdateConstrainCuda::Impl::getCoordinatesReadySync()
+GpuEventSynchronizer* UpdateConstrainGpu::Impl::getCoordinatesReadySync()
 {
     return coordinatesReady_;
 }
 
-UpdateConstrainCuda::UpdateConstrainCuda(const t_inputrec&     ir,
-                                         const gmx_mtop_t&     mtop,
-                                         const void*           commandStream,
-                                         GpuEventSynchronizer* xUpdatedOnDevice) :
+UpdateConstrainGpu::UpdateConstrainGpu(const t_inputrec&     ir,
+                                       const gmx_mtop_t&     mtop,
+                                       const void*           commandStream,
+                                       GpuEventSynchronizer* xUpdatedOnDevice) :
     impl_(new Impl(ir, mtop, commandStream, xUpdatedOnDevice))
 {
 }
 
-UpdateConstrainCuda::~UpdateConstrainCuda() = default;
-
-void UpdateConstrainCuda::integrate(GpuEventSynchronizer*             fReadyOnDevice,
-                                    const real                        dt,
-                                    const bool                        updateVelocities,
-                                    const bool                        computeVirial,
-                                    tensor                            virialScaled,
-                                    const bool                        doTemperatureScaling,
-                                    gmx::ArrayRef<const t_grp_tcstat> tcstat,
-                                    const bool                        doParrinelloRahman,
-                                    const float                       dtPressureCouple,
-                                    const matrix                      prVelocityScalingMatrix)
+UpdateConstrainGpu::~UpdateConstrainGpu() = default;
+
+void UpdateConstrainGpu::integrate(GpuEventSynchronizer*             fReadyOnDevice,
+                                   const real                        dt,
+                                   const bool                        updateVelocities,
+                                   const bool                        computeVirial,
+                                   tensor                            virialScaled,
+                                   const bool                        doTemperatureScaling,
+                                   gmx::ArrayRef<const t_grp_tcstat> tcstat,
+                                   const bool                        doParrinelloRahman,
+                                   const float                       dtPressureCouple,
+                                   const matrix                      prVelocityScalingMatrix)
 {
     impl_->integrate(fReadyOnDevice, dt, updateVelocities, computeVirial, virialScaled, doTemperatureScaling,
                      tcstat, doParrinelloRahman, dtPressureCouple, prVelocityScalingMatrix);
 }
 
-void UpdateConstrainCuda::scaleCoordinates(const matrix scalingMatrix)
+void UpdateConstrainGpu::scaleCoordinates(const matrix scalingMatrix)
 {
     impl_->scaleCoordinates(scalingMatrix);
 }
 
-void UpdateConstrainCuda::set(DeviceBuffer<float>       d_x,
-                              DeviceBuffer<float>       d_v,
-                              const DeviceBuffer<float> d_f,
-                              const t_idef&             idef,
-                              const t_mdatoms&          md,
-                              const int                 numTempScaleValues)
+void UpdateConstrainGpu::set(DeviceBuffer<float>       d_x,
+                             DeviceBuffer<float>       d_v,
+                             const DeviceBuffer<float> d_f,
+                             const t_idef&             idef,
+                             const t_mdatoms&          md,
+                             const int                 numTempScaleValues)
 {
     impl_->set(d_x, d_v, d_f, idef, md, numTempScaleValues);
 }
 
-void UpdateConstrainCuda::setPbc(const PbcType pbcType, const matrix box)
+void UpdateConstrainGpu::setPbc(const PbcType pbcType, const matrix box)
 {
     impl_->setPbc(pbcType, box);
 }
 
-GpuEventSynchronizer* UpdateConstrainCuda::getCoordinatesReadySync()
+GpuEventSynchronizer* UpdateConstrainGpu::getCoordinatesReadySync()
 {
     return impl_->getCoordinatesReadySync();
 }
 
-bool UpdateConstrainCuda::isNumCoupledConstraintsSupported(const gmx_mtop_t& mtop)
+bool UpdateConstrainGpu::isNumCoupledConstraintsSupported(const gmx_mtop_t& mtop)
 {
     return LincsGpu::isNumCoupledConstraintsSupported(mtop);
 }

diff --git a/src/gromacs/mdlib/update_constrain_cuda_impl.h b/src/gromacs/mdlib/update_constrain_gpu_impl.h
similarity index 93%
rename from src/gromacs/mdlib/update_constrain_cuda_impl.h
rename to src/gromacs/mdlib/update_constrain_gpu_impl.h
index dcb48f9ca7e5b21aaf11947421022c01b6848401..5a28045afd48e82b21ec62cb91864e06233ed81e 100644 (file)
--- a/src/gromacs/mdlib/update_constrain_cuda_impl.h
+++ b/src/gromacs/mdlib/update_constrain_gpu_impl.h
@@ -34,7 +34,7 @@
  */
 /*! \internal \file
  *
- * \brief Declares CUDA implementation class for update and constraints.
+ * \brief Declares GPU implementation class for update and constraints.
  *
  * This header file is needed to include from both the device-side
  * kernels file, and the host-side management code.
@@ -43,8 +43,8 @@
  *
  * \ingroup module_mdlib
  */
-#ifndef GMX_MDLIB_UPDATE_CONSTRAIN_CUDA_IMPL_H
-#define GMX_MDLIB_UPDATE_CONSTRAIN_CUDA_IMPL_H
+#ifndef GMX_MDLIB_UPDATE_CONSTRAIN_GPU_IMPL_H
+#define GMX_MDLIB_UPDATE_CONSTRAIN_GPU_IMPL_H
 
 #include "gmxpre.h"
 
@@ -52,14 +52,14 @@
 #include "gromacs/mdlib/leapfrog_gpu.cuh"
 #include "gromacs/mdlib/lincs_gpu.cuh"
 #include "gromacs/mdlib/settle_gpu.cuh"
-#include "gromacs/mdlib/update_constrain_cuda.h"
+#include "gromacs/mdlib/update_constrain_gpu.h"
 #include "gromacs/mdtypes/inputrec.h"
 
 namespace gmx
 {
 
-/*! \internal \brief Class with interfaces and data for CUDA version of Update-Constraint. */
-class UpdateConstrainCuda::Impl
+/*! \internal \brief Class with interfaces and data for GPU version of Update-Constraint. */
+class UpdateConstrainGpu::Impl
 {
 
 public:
@@ -156,16 +156,16 @@ public:
 
     /*! \brief
      * Returns whether the maximum number of coupled constraints is supported
-     * by the CUDA LINCS code.
+     * by the GPU LINCS code.
      *
      * \param[in] mtop The molecular topology
      */
     static bool isNumCoupledConstraintsSupported(const gmx_mtop_t& mtop);
 
 private:
-    //! CUDA stream
+    //! GPU stream
     CommandStream commandStream_ = nullptr;
-    //! CUDA kernel launch config
+    //! GPU kernel launch config
     KernelLaunchConfig coordinateScalingKernelLaunchConfig_;
 
     //! Periodic boundary data
@@ -198,9 +198,9 @@ private:
 
     //! Leap-Frog integrator
     std::unique_ptr<LeapFrogGpu> integrator_;
-    //! LINCS CUDA object to use for non-water constraints
+    //! LINCS GPU object to use for non-water constraints
     std::unique_ptr<LincsGpu> lincsGpu_;
-    //! SETTLE CUDA object for water constrains
+    //! SETTLE GPU object for water constrains
     std::unique_ptr<SettleGpu> settleGpu_;
 
     //! An pointer to the event to indicate when the update of coordinates is complete
@@ -209,4 +209,4 @@ private:
 
 } // namespace gmx
 
-#endif
+#endif // GMX_MDLIB_UPDATE_CONSTRAIN_GPU_IMPL_H

diff --git a/src/gromacs/mdrun/md.cpp b/src/gromacs/mdrun/md.cpp
index 68dc52d2d6fa7f22023c215a267a726264fa8b91..91d932e3d8fd286f28cb7529523cc371f19e1cbc 100644 (file)
--- a/src/gromacs/mdrun/md.cpp
+++ b/src/gromacs/mdrun/md.cpp
@@ -95,7 +95,7 @@
 #include "gromacs/mdlib/tgroup.h"
 #include "gromacs/mdlib/trajectory_writing.h"
 #include "gromacs/mdlib/update.h"
-#include "gromacs/mdlib/update_constrain_cuda.h"
+#include "gromacs/mdlib/update_constrain_gpu.h"
 #include "gromacs/mdlib/vcm.h"
 #include "gromacs/mdlib/vsite.h"
 #include "gromacs/mdrunutility/handlerestart.h"
@@ -287,7 +287,7 @@ void gmx::LegacySimulator::do_md()
 
     auto mdatoms = mdAtoms->mdatoms();
 
-    std::unique_ptr<UpdateConstrainCuda> integrator;
+    std::unique_ptr<UpdateConstrainGpu> integrator;
 
     if (DOMAINDECOMP(cr))
     {
@@ -367,7 +367,7 @@ void gmx::LegacySimulator::do_md()
         {
             GMX_LOG(mdlog.info).asParagraph().appendText("Updating coordinates on the GPU.");
         }
-        integrator = std::make_unique<UpdateConstrainCuda>(
+        integrator = std::make_unique<UpdateConstrainGpu>(
                 *ir, *top_global, stateGpu->getUpdateStream(), stateGpu->xUpdatedOnDevice());
 
         integrator->setPbc(PbcType::Xyz, state->box);

diff --git a/src/gromacs/taskassignment/decidegpuusage.cpp b/src/gromacs/taskassignment/decidegpuusage.cpp
index 6df6e2322b531e8e7cce81d01a1438726dcc6020..68bc2c64abbfc446cfb5696cf9cec1d49dd1b096 100644 (file)
--- a/src/gromacs/taskassignment/decidegpuusage.cpp
+++ b/src/gromacs/taskassignment/decidegpuusage.cpp
@@ -58,7 +58,7 @@
 #include "gromacs/hardware/hardwaretopology.h"
 #include "gromacs/hardware/hw_info.h"
 #include "gromacs/mdlib/gmx_omp_nthreads.h"
-#include "gromacs/mdlib/update_constrain_cuda.h"
+#include "gromacs/mdlib/update_constrain_gpu.h"
 #include "gromacs/mdtypes/commrec.h"
 #include "gromacs/mdtypes/inputrec.h"
 #include "gromacs/mdtypes/md_enums.h"
@@ -615,10 +615,10 @@ bool decideWhetherToUseGpuForUpdate(const bool           forceGpuUpdateDefault,
     {
         errorMessage += "Non-connecting constraints are not supported";
     }
-    if (!UpdateConstrainCuda::isNumCoupledConstraintsSupported(mtop))
+    if (!UpdateConstrainGpu::isNumCoupledConstraintsSupported(mtop))
     {
         errorMessage +=
-                "The number of coupled constraints is higher than supported in the CUDA LINCS "
+                "The number of coupled constraints is higher than supported in the GPU LINCS "
                 "code.\n";
     }