Fix clang-tidy warnings in the OpenCL build

Change-Id: I4cc6a1402d068431feab4523911a7b60e0e66d59

diff --git a/src/gromacs/ewald/pme_gpu_internal.cpp b/src/gromacs/ewald/pme_gpu_internal.cpp
index 70c74c900ce146a35949cde6483407790e22b472..e99e7a9e9332f726fb37f97da3c997427dd73651 100644 (file)
--- a/src/gromacs/ewald/pme_gpu_internal.cpp
+++ b/src/gromacs/ewald/pme_gpu_internal.cpp
@@ -642,13 +642,13 @@ PmeOutput pme_gpu_getEnergyAndVirial(const gmx_pme_t &pme)
     }
 
     unsigned int j = 0;
-    output.coulombVirial_[XX][XX] = 0.25f * pmeGpu->staging.h_virialAndEnergy[j++];
-    output.coulombVirial_[YY][YY] = 0.25f * pmeGpu->staging.h_virialAndEnergy[j++];
-    output.coulombVirial_[ZZ][ZZ] = 0.25f * pmeGpu->staging.h_virialAndEnergy[j++];
-    output.coulombVirial_[XX][YY] = output.coulombVirial_[YY][XX] = 0.25f * pmeGpu->staging.h_virialAndEnergy[j++];
-    output.coulombVirial_[XX][ZZ] = output.coulombVirial_[ZZ][XX] = 0.25f * pmeGpu->staging.h_virialAndEnergy[j++];
-    output.coulombVirial_[YY][ZZ] = output.coulombVirial_[ZZ][YY] = 0.25f * pmeGpu->staging.h_virialAndEnergy[j++];
-    output.coulombEnergy_         = 0.5f * pmeGpu->staging.h_virialAndEnergy[j++];
+    output.coulombVirial_[XX][XX] = 0.25F * pmeGpu->staging.h_virialAndEnergy[j++];
+    output.coulombVirial_[YY][YY] = 0.25F * pmeGpu->staging.h_virialAndEnergy[j++];
+    output.coulombVirial_[ZZ][ZZ] = 0.25F * pmeGpu->staging.h_virialAndEnergy[j++];
+    output.coulombVirial_[XX][YY] = output.coulombVirial_[YY][XX] = 0.25F * pmeGpu->staging.h_virialAndEnergy[j++];
+    output.coulombVirial_[XX][ZZ] = output.coulombVirial_[ZZ][XX] = 0.25F * pmeGpu->staging.h_virialAndEnergy[j++];
+    output.coulombVirial_[YY][ZZ] = output.coulombVirial_[ZZ][YY] = 0.25F * pmeGpu->staging.h_virialAndEnergy[j++];
+    output.coulombEnergy_         = 0.5F * pmeGpu->staging.h_virialAndEnergy[j++];
 
     return output;
 }

diff --git a/src/gromacs/ewald/pme_gpu_program_impl_ocl.cpp b/src/gromacs/ewald/pme_gpu_program_impl_ocl.cpp
index a3558005608581aa08af5fc05eaa6b57aa517ea2..492128ed3eeaa9e90d852893ff6cc0dbf50367db 100644 (file)
--- a/src/gromacs/ewald/pme_gpu_program_impl_ocl.cpp
+++ b/src/gromacs/ewald/pme_gpu_program_impl_ocl.cpp
@@ -110,7 +110,7 @@ PmeGpuProgramImpl::~PmeGpuProgramImpl()
  * smaller than the minimum order^2 required in spread/gather ATM which
  * we need to check for.
  */
-static void checkRequiredWarpSize(const cl_kernel          kernel,
+static void checkRequiredWarpSize(cl_kernel                kernel,
                                   const char*              kernelName,
                                   const gmx_device_info_t *deviceInfo)
 {
@@ -164,8 +164,8 @@ void PmeGpuProgramImpl::compileKernels(const gmx_device_info_t *deviceInfo)
                     warpSize / c_pmeSpreadGatherThreadsPerAtom,
                     c_pmeSpreadGatherThreadsPerAtom,
                     static_cast<float>(c_pmeMaxUnitcellShift),
-                    c_usePadding,
-                    c_skipNeutralAtoms,
+                    static_cast<int>(c_usePadding),
+                    static_cast<int>(c_skipNeutralAtoms),
                     c_virialAndEnergyCount,
                     spreadWorkGroupSize,
                     solveMaxWorkGroupSize,

diff --git a/src/gromacs/gpu_utils/gpu_utils_ocl.cpp b/src/gromacs/gpu_utils/gpu_utils_ocl.cpp
index 2dd684891637e7da81c4f660b07a170148fff0e7..2ddde385f5fc842a0a464c4ac6cc86d153409070 100644 (file)
--- a/src/gromacs/gpu_utils/gpu_utils_ocl.cpp
+++ b/src/gromacs/gpu_utils/gpu_utils_ocl.cpp
@@ -149,7 +149,7 @@ static bool isDeviceSane(const gmx_device_info_t *devInfo,
 {
     cl_context_properties properties[] = {
         CL_CONTEXT_PLATFORM,
-        (cl_context_properties) devInfo->ocl_gpu_id.ocl_platform_id,
+        reinterpret_cast<cl_context_properties>(devInfo->ocl_gpu_id.ocl_platform_id),
         0
     };
     // uncrustify spacing
@@ -286,7 +286,7 @@ static int checkGpu(size_t                   deviceId,
     return egpuCompatible;
 }
 
-} // namespace
+} // namespace gmx
 
 /*! \brief Returns an ocl_vendor_id_t value corresponding to the input OpenCL vendor name.
  *

diff --git a/src/gromacs/gpu_utils/oclraii.h b/src/gromacs/gpu_utils/oclraii.h
index 227ba844e536cd541ff12cfaf993538aa4c409f4..23c77ef14ee16aba53104083ba3c5f6e5dd31f95 100644 (file)
--- a/src/gromacs/gpu_utils/oclraii.h
+++ b/src/gromacs/gpu_utils/oclraii.h
@@ -114,9 +114,9 @@ class ClHandle
         //! Deleted copy constructor.
         ClHandle(const ClHandle &)            = delete;
         //! Default move assignment operator.
-        ClHandle &operator=(ClHandle &&)      = default;
+        ClHandle &operator=(ClHandle &&) noexcept = default;
         //! Default copy constructor.
-        ClHandle(ClHandle &&)                 = default;
+        ClHandle(ClHandle &&)  noexcept               = default;
         /*! \brief Convenience conversion operator so the wrapper type
          * can simply convert to the wrapped type. */
         operator cl_type() const { return handle_; }
@@ -133,6 +133,6 @@ using ClProgram      = ClHandle<cl_program>;
 using ClKernel       = ClHandle<cl_kernel>;
 /*! @} */
 
-} // namespace
+} // namespace gmx
 
 #endif

diff --git a/src/gromacs/gpu_utils/tests/devicetransfers_ocl.cpp b/src/gromacs/gpu_utils/tests/devicetransfers_ocl.cpp
index 7c0215ef26ee92479f1a66b90309dbce9e8e0949..8cb3e99aac9185235542db74adefcc83e6469aa8 100644 (file)
--- a/src/gromacs/gpu_utils/tests/devicetransfers_ocl.cpp
+++ b/src/gromacs/gpu_utils/tests/devicetransfers_ocl.cpp
@@ -1,7 +1,7 @@
 /*
  * This file is part of the GROMACS molecular simulation package.
  *
- * Copyright (c) 2017,2018, by the GROMACS development team, led by
+ * Copyright (c) 2017,2018,2019, 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.
@@ -60,7 +60,7 @@ namespace
  *
  * \throws InternalError  If status indicates failure, supplying
  *                        descriptive text from \c message. */
-static void throwUponFailure(cl_int status, const char *message)
+void throwUponFailure(cl_int status, const char *message)
 {
     if (status != CL_SUCCESS)
     {
@@ -86,13 +86,13 @@ void doDeviceTransfers(const gmx_gpu_info_t &gpuInfo,
     const auto           *device       = getDeviceInfo(gpuInfo, compatibleGpus[0]);
     cl_context_properties properties[] = {
         CL_CONTEXT_PLATFORM,
-        (cl_context_properties) device->ocl_gpu_id.ocl_platform_id,
+        reinterpret_cast<cl_context_properties>(device->ocl_gpu_id.ocl_platform_id),
         0
     };
     // Give uncrustify more space
 
     auto deviceId = device->ocl_gpu_id.ocl_device_id;
-    auto context  = clCreateContext(properties, 1, &deviceId, NULL, NULL, &status);
+    auto context  = clCreateContext(properties, 1, &deviceId, nullptr, nullptr, &status);
     throwUponFailure(status, "creating context");
     auto commandQueue = clCreateCommandQueue(context, deviceId, 0, &status);
     throwUponFailure(status, "creating command queue");

diff --git a/src/gromacs/mdlib/update.cpp b/src/gromacs/mdlib/update.cpp
index 2e76a401ddc3d4b00b3ba49b130a6da2017389a5..adb0694f7dcb3e86cc035e57f9069233fa66b5f6 100644 (file)
--- a/src/gromacs/mdlib/update.cpp
+++ b/src/gromacs/mdlib/update.cpp
@@ -1654,11 +1654,11 @@ void finish_update(const t_inputrec       *inputrec, /* input record and box stu
         }
         else
         {
-            auto           xp = makeConstArrayRef(*upd->xp()).subArray(0, homenr);
-            auto           x  = makeArrayRef(state->x).subArray(0, homenr);
-#ifndef __clang_analyzer__
+            auto xp  = makeConstArrayRef(*upd->xp()).subArray(0, homenr);
+            auto x   = makeArrayRef(state->x).subArray(0, homenr);
+
+
             int gmx_unused nth = gmx_omp_nthreads_get(emntUpdate);
-#endif
 #pragma omp parallel for num_threads(nth) schedule(static)
             for (int i = 0; i < homenr; i++)
             {

diff --git a/src/gromacs/nbnxm/gpu_common.h b/src/gromacs/nbnxm/gpu_common.h
index fc5977c7c38d9f88e08877a791a46d399863719a..151dc4b3ff6fac2ef608762e41d153fec89b8443 100644 (file)
--- a/src/gromacs/nbnxm/gpu_common.h
+++ b/src/gromacs/nbnxm/gpu_common.h
@@ -124,10 +124,10 @@ gpuAtomToInteractionLocality(const AtomLocality atomLocality)
 }
 
 
-void
-setupGpuShortRangeWork(gmx_nbnxn_gpu_t                  *nb,
-                       const gmx::GpuBonded             *gpuBonded,
-                       const Nbnxm::InteractionLocality  iLocality)
+//NOLINTNEXTLINE(misc-definitions-in-headers)
+void setupGpuShortRangeWork(gmx_nbnxn_gpu_t                  *nb,
+                            const gmx::GpuBonded             *gpuBonded,
+                            const Nbnxm::InteractionLocality  iLocality)
 {
     GMX_ASSERT(nb, "Need a valid nbnxn_gpu object");
 
@@ -155,9 +155,9 @@ haveGpuShortRangeWork(const gmx_nbnxn_gpu_t            &nb,
     return nb.haveWork[iLocality];
 }
 
-bool
-haveGpuShortRangeWork(const gmx_nbnxn_gpu_t     *nb,
-                      const Nbnxm::AtomLocality  aLocality)
+//NOLINTNEXTLINE(misc-definitions-in-headers)
+bool haveGpuShortRangeWork(const gmx_nbnxn_gpu_t     *nb,
+                           const Nbnxm::AtomLocality  aLocality)
 {
     GMX_ASSERT(nb, "Need a valid nbnxn_gpu object");
 

diff --git a/src/gromacs/nbnxm/nbnxm_gpu.h b/src/gromacs/nbnxm/nbnxm_gpu.h
index ba3b9a8f52778b067baf6fd2ce3e7cf59310beef..93d608347b830c21a02cd0fd51266ceed14c5747 100644 (file)
--- a/src/gromacs/nbnxm/nbnxm_gpu.h
+++ b/src/gromacs/nbnxm/nbnxm_gpu.h
@@ -144,7 +144,7 @@ void gpu_launch_cpyback(gmx_nbnxn_gpu_t  gmx_unused *nb,
                         nbnxn_atomdata_t gmx_unused *nbatom,
                         int              gmx_unused  flags,
                         AtomLocality     gmx_unused  aloc,
-                        const bool       gmx_unused  copyBackNbForce) GPU_FUNC_TERM;
+                        bool             gmx_unused  copyBackNbForce) GPU_FUNC_TERM;
 
 /*! \brief Attempts to complete nonbonded GPU task.
  *
@@ -254,7 +254,7 @@ void nbnxnInsertNonlocalGpuDependency(const gmx_nbnxn_gpu_t gmx_unused    *nb,
 GPU_FUNC_QUALIFIER
 void setupGpuShortRangeWork(gmx_nbnxn_gpu_t                  gmx_unused *nb,
                             const gmx::GpuBonded             gmx_unused *gpuBonded,
-                            const Nbnxm::InteractionLocality gmx_unused  iLocality) GPU_FUNC_TERM;
+                            Nbnxm::InteractionLocality       gmx_unused  iLocality) GPU_FUNC_TERM;
 
 /*! \brief Returns true if there is GPU short-range work for the given atom locality.
  *
@@ -267,7 +267,7 @@ void setupGpuShortRangeWork(gmx_nbnxn_gpu_t                  gmx_unused *nb,
  */
 GPU_FUNC_QUALIFIER
 bool haveGpuShortRangeWork(const gmx_nbnxn_gpu_t     gmx_unused *nb,
-                           const Nbnxm::AtomLocality gmx_unused  aLocality) GPU_FUNC_TERM_WITH_RETURN(false);
+                           Nbnxm::AtomLocality       gmx_unused  aLocality) GPU_FUNC_TERM_WITH_RETURN(false);
 
 /*! \brief Initialization for F buffer operations on GPU */
 CUDA_FUNC_QUALIFIER

diff --git a/src/gromacs/nbnxm/opencl/nbnxm_ocl.cpp b/src/gromacs/nbnxm/opencl/nbnxm_ocl.cpp
index 3ac8491795893aaad002e95cecd44a7434c0a875..99a7829e29df983e782571f6fab14c6f9e339ba9 100644 (file)
--- a/src/gromacs/nbnxm/opencl/nbnxm_ocl.cpp
+++ b/src/gromacs/nbnxm/opencl/nbnxm_ocl.cpp
@@ -133,7 +133,7 @@ static inline void validate_global_work_size(const KernelLaunchConfig &config, i
     {
         size_t device_limit;
 
-        device_limit = (1ull << device_size_t_size_bits) - 1;
+        device_limit = (1ULL << device_size_t_size_bits) - 1;
 
         for (int i = 0; i < work_dim; i++)
         {

diff --git a/src/gromacs/nbnxm/opencl/nbnxm_ocl_data_mgmt.cpp b/src/gromacs/nbnxm/opencl/nbnxm_ocl_data_mgmt.cpp
index 1f00c06808e62a92b22e9fb07b54ad5faa37b185..843ffc04331213704fef25a103bfba0250af096c 100644 (file)
--- a/src/gromacs/nbnxm/opencl/nbnxm_ocl_data_mgmt.cpp
+++ b/src/gromacs/nbnxm/opencl/nbnxm_ocl_data_mgmt.cpp
@@ -488,7 +488,7 @@ static void init_timings(gmx_wallclock_gpu_nbnxn_t *t)
 
 //! OpenCL notification callback function
 static void CL_CALLBACK
-ocl_notify_fn( const char *pErrInfo, const void *, size_t, void *)
+ocl_notify_fn(const char *pErrInfo, const void gmx_unused *private_info, size_t gmx_unused cb, void gmx_unused *user_data)
 {
     if (pErrInfo != nullptr)
     {
@@ -762,7 +762,7 @@ static void nbnxn_ocl_clear_f(gmx_nbnxn_ocl_t *nb, int natoms_clear)
 
     cl_atomdata_t           *atomData = nb->atdat;
     cl_command_queue         ls       = nb->stream[InteractionLocality::Local];
-    cl_float                 value    = 0.0f;
+    cl_float                 value    = 0.0F;
 
     cl_error = clEnqueueFillBuffer(ls, atomData->f, &value, sizeof(cl_float),
                                    0, natoms_clear*sizeof(rvec), 0, nullptr, nullptr);

diff --git a/src/gromacs/nbnxm/pairlist.cpp b/src/gromacs/nbnxm/pairlist.cpp
index 296db78c92d7a8ff308b4d73f0ba1309f4ed100a..6da7cda8145bc74603772a5566a2e0c025451a7d 100644 (file)
--- a/src/gromacs/nbnxm/pairlist.cpp
+++ b/src/gromacs/nbnxm/pairlist.cpp
@@ -2693,7 +2693,7 @@ static void combine_nblists(gmx::ArrayRef<const NbnxnPairlistGpu>  nbls,
     /* Each thread should copy its own data to the combined arrays,
      * as otherwise data will go back and forth between different caches.
      */
-    int gmx_unused nthreads = gmx_omp_nthreads_get(emntPairsearch);
+    const int gmx_unused nthreads = gmx_omp_nthreads_get(emntPairsearch);
 
 #pragma omp parallel for num_threads(nthreads) schedule(static)
     for (gmx::index n = 0; n < nbls.ssize(); n++)