From bf9b4cd5aa623e089a7432d38e47bee27103a478 Mon Sep 17 00:00:00 2001 From: =?utf8?q?Szil=C3=A1rd=20P=C3=A1ll?= Date: Wed, 21 Aug 2019 14:02:15 +0200 Subject: [PATCH] Fix clang-tidy warnings in the OpenCL build Change-Id: I4cc6a1402d068431feab4523911a7b60e0e66d59 --- src/gromacs/ewald/pme_gpu_internal.cpp | 14 +++++++------- src/gromacs/ewald/pme_gpu_program_impl_ocl.cpp | 6 +++--- src/gromacs/gpu_utils/gpu_utils_ocl.cpp | 4 ++-- src/gromacs/gpu_utils/oclraii.h | 6 +++--- .../gpu_utils/tests/devicetransfers_ocl.cpp | 8 ++++---- src/gromacs/mdlib/update.cpp | 8 ++++---- src/gromacs/nbnxm/gpu_common.h | 14 +++++++------- src/gromacs/nbnxm/nbnxm_gpu.h | 6 +++--- src/gromacs/nbnxm/opencl/nbnxm_ocl.cpp | 2 +- src/gromacs/nbnxm/opencl/nbnxm_ocl_data_mgmt.cpp | 4 ++-- src/gromacs/nbnxm/pairlist.cpp | 2 +- 11 files changed, 37 insertions(+), 37 deletions(-) diff --git a/src/gromacs/ewald/pme_gpu_internal.cpp b/src/gromacs/ewald/pme_gpu_internal.cpp index 70c74c900c..e99e7a9e93 100644 --- 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 a355800560..492128ed3e 100644 --- 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(c_pmeMaxUnitcellShift), - c_usePadding, - c_skipNeutralAtoms, + static_cast(c_usePadding), + static_cast(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 2dd6848916..2ddde385f5 100644 --- 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(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 227ba844e5..23c77ef14e 100644 --- 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; using ClKernel = ClHandle; /*! @} */ -} // 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 7c0215ef26..8cb3e99aac 100644 --- 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(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 2e76a401dd..adb0694f7d 100644 --- 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 fc5977c7c3..151dc4b3ff 100644 --- 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 ba3b9a8f52..93d608347b 100644 --- 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 3ac8491795..99a7829e29 100644 --- 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 1f00c06808..843ffc0433 100644 --- 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 296db78c92..6da7cda814 100644 --- a/src/gromacs/nbnxm/pairlist.cpp +++ b/src/gromacs/nbnxm/pairlist.cpp @@ -2693,7 +2693,7 @@ static void combine_nblists(gmx::ArrayRef 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++) -- 2.22.0