From b051ca85dbb7d07d247798f6246b345f4456f05e Mon Sep 17 00:00:00 2001 From: Andrey Alekseenko Date: Fri, 12 Feb 2021 15:42:10 +0000 Subject: [PATCH] Add gmx::isPowerOfTwo function A follow-up from a discussion in !1080. --- api/legacy/include/gromacs/math/functions.h | 16 ++++++++- src/gromacs/ewald/pme_gather.cu | 5 +-- src/gromacs/ewald/pme_solve.cpp | 5 ++- src/gromacs/math/tests/functions.cpp | 36 +++++++++++++++++++- src/gromacs/mdlib/lincs_gpu.cu | 5 +-- src/gromacs/mdlib/settle_gpu.cu | 5 +-- src/gromacs/nbnxm/benchmark/bench_system.cpp | 4 +-- src/gromacs/nbnxm/nbnxm_geometry.h | 10 ++---- 8 files changed, 65 insertions(+), 21 deletions(-) diff --git a/api/legacy/include/gromacs/math/functions.h b/api/legacy/include/gromacs/math/functions.h index 271e1cf760..af557fb6ad 100644 --- a/api/legacy/include/gromacs/math/functions.h +++ b/api/legacy/include/gromacs/math/functions.h @@ -1,7 +1,8 @@ /* * This file is part of the GROMACS molecular simulation package. * - * Copyright (c) 2015,2016,2018,2019,2020, by the GROMACS development team, led by + * Copyright (c) 2015,2016,2018,2019,2020 by the GROMACS development team. + * Copyright (c) 2021, 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. @@ -472,6 +473,19 @@ static inline int64_t roundToInt64(double x) return static_cast(rint(x)); } +//! \brief Check whether \p v is an integer power of 2. +template::value>> +#if defined(__NVCC__) && !defined(__CUDACC_RELAXED_CONSTEXPR__) +/* In CUDA 11, a constexpr function cannot be called from a function with incompatible execution + * space, unless --expt-relaxed-constexpr flag is set */ +__host__ __device__ +#endif + static inline constexpr bool + isPowerOfTwo(const T v) +{ + return (v > 0) && ((v & (v - 1)) == 0); +} + } // namespace gmx diff --git a/src/gromacs/ewald/pme_gather.cu b/src/gromacs/ewald/pme_gather.cu index a81791765c..35111b52ec 100644 --- a/src/gromacs/ewald/pme_gather.cu +++ b/src/gromacs/ewald/pme_gather.cu @@ -1,7 +1,8 @@ /* * This file is part of the GROMACS molecular simulation package. * - * Copyright (c) 2016,2017,2018,2019,2020, by the GROMACS development team, led by + * Copyright (c) 2016,2017,2018,2019,2020 by the GROMACS development team. + * Copyright (c) 2021, 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. @@ -92,7 +93,7 @@ __device__ __forceinline__ void reduce_atom_forces(float3* __restrict__ sm_force float& fy, float& fz) { - if (!(order & (order - 1))) // Only for orders of power of 2 + if (gmx::isPowerOfTwo(order)) // Only for orders of power of 2 { const unsigned int activeMask = c_fullWarpMask; diff --git a/src/gromacs/ewald/pme_solve.cpp b/src/gromacs/ewald/pme_solve.cpp index 88e529b553..03ca35bfee 100644 --- a/src/gromacs/ewald/pme_solve.cpp +++ b/src/gromacs/ewald/pme_solve.cpp @@ -4,7 +4,7 @@ * Copyright (c) 1991-2000, University of Groningen, The Netherlands. * Copyright (c) 2001-2004, The GROMACS development team. * Copyright (c) 2013,2014,2015,2016,2017 by the GROMACS development team. - * Copyright (c) 2018,2019,2020, by the GROMACS development team, led by + * Copyright (c) 2018,2019,2020,2021, 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. @@ -94,8 +94,7 @@ constexpr int c_simdWidth = 4; template static size_t roundUpToMultipleOfFactor(size_t number) { - static_assert(factor > 0 && (factor & (factor - 1)) == 0, - "factor should be >0 and a power of 2"); + static_assert(gmx::isPowerOfTwo(factor)); /* We need to add a most factor-1 and because factor is a power of 2, * we get the result by masking out the bits corresponding to factor-1. diff --git a/src/gromacs/math/tests/functions.cpp b/src/gromacs/math/tests/functions.cpp index b444ee7841..939a08da2e 100644 --- a/src/gromacs/math/tests/functions.cpp +++ b/src/gromacs/math/tests/functions.cpp @@ -1,7 +1,8 @@ /* * This file is part of the GROMACS molecular simulation package. * - * Copyright (c) 2015,2016,2017,2018,2019, by the GROMACS development team, led by + * Copyright (c) 2015,2016,2017,2018,2019 by the GROMACS development team. + * Copyright (c) 2021, 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. @@ -356,4 +357,37 @@ TEST(FunctionTest, ErfAndErfInvAreInversesDouble) } } +template +class FunctionTestIntegerTypes : public ::testing::Test +{ +}; + +typedef ::testing::Types IntegerTypes; +TYPED_TEST_CASE(FunctionTestIntegerTypes, IntegerTypes); + +TYPED_TEST(FunctionTestIntegerTypes, IsPowerOfTwo) +{ + if (std::is_signed_v) + { + EXPECT_EQ(false, gmx::isPowerOfTwo(std::numeric_limits::min())); + EXPECT_EQ(false, gmx::isPowerOfTwo(-16)); + EXPECT_EQ(false, gmx::isPowerOfTwo(-3)); + EXPECT_EQ(false, gmx::isPowerOfTwo(-2)); + EXPECT_EQ(false, gmx::isPowerOfTwo(-1)); + } + EXPECT_EQ(false, gmx::isPowerOfTwo(0)); + EXPECT_EQ(true, gmx::isPowerOfTwo(1)); + EXPECT_EQ(true, gmx::isPowerOfTwo(2)); + EXPECT_EQ(false, gmx::isPowerOfTwo(3)); + EXPECT_EQ(true, gmx::isPowerOfTwo(4)); + EXPECT_EQ(false, gmx::isPowerOfTwo(5)); + EXPECT_EQ(false, gmx::isPowerOfTwo(6)); + EXPECT_EQ(false, gmx::isPowerOfTwo(24)); + EXPECT_EQ(false, gmx::isPowerOfTwo(63)); + EXPECT_EQ(true, gmx::isPowerOfTwo(64)); + EXPECT_EQ(false, gmx::isPowerOfTwo(66)); + // Max for any type is always 2^x - 1 + EXPECT_EQ(false, gmx::isPowerOfTwo(std::numeric_limits::max())); +} + } // namespace diff --git a/src/gromacs/mdlib/lincs_gpu.cu b/src/gromacs/mdlib/lincs_gpu.cu index ac879bcfae..0967c20781 100644 --- a/src/gromacs/mdlib/lincs_gpu.cu +++ b/src/gromacs/mdlib/lincs_gpu.cu @@ -1,7 +1,7 @@ /* * This file is part of the GROMACS molecular simulation package. * - * Copyright (c) 2019,2020, by the GROMACS development team, led by + * Copyright (c) 2019,2020,2021, 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. @@ -61,6 +61,7 @@ #include "gromacs/gpu_utils/devicebuffer.cuh" #include "gromacs/gpu_utils/gputraits.cuh" #include "gromacs/gpu_utils/vectype_ops.cuh" +#include "gromacs/math/functions.h" #include "gromacs/math/vec.h" #include "gromacs/mdlib/constr.h" #include "gromacs/pbcutil/pbc.h" @@ -529,7 +530,7 @@ LincsGpu::LincsGpu(int numIterations, static_assert(sizeof(real) == sizeof(float), "Real numbers should be in single precision in GPU code."); static_assert( - c_threadsPerBlock > 0 && ((c_threadsPerBlock & (c_threadsPerBlock - 1)) == 0), + gmx::isPowerOfTwo(c_threadsPerBlock), "Number of threads per block should be a power of two in order for reduction to work."); allocateDeviceBuffer(&kernelParams_.d_virialScaled, 6, deviceContext_); diff --git a/src/gromacs/mdlib/settle_gpu.cu b/src/gromacs/mdlib/settle_gpu.cu index eb4f7b066b..eb4b495772 100644 --- a/src/gromacs/mdlib/settle_gpu.cu +++ b/src/gromacs/mdlib/settle_gpu.cu @@ -1,7 +1,7 @@ /* * This file is part of the GROMACS molecular simulation package. * - * Copyright (c) 2019,2020, by the GROMACS development team, led by + * Copyright (c) 2019,2020,2021, 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. @@ -61,6 +61,7 @@ #include "gromacs/gpu_utils/devicebuffer.h" #include "gromacs/gpu_utils/gputraits.cuh" #include "gromacs/gpu_utils/vectype_ops.cuh" +#include "gromacs/math/functions.h" #include "gromacs/math/vec.h" #include "gromacs/pbcutil/pbc.h" #include "gromacs/pbcutil/pbc_aiuc_cuda.cuh" @@ -478,7 +479,7 @@ SettleGpu::SettleGpu(const gmx_mtop_t& mtop, const DeviceContext& deviceContext, static_assert(sizeof(real) == sizeof(float), "Real numbers should be in single precision in GPU code."); static_assert( - c_threadsPerBlock > 0 && ((c_threadsPerBlock & (c_threadsPerBlock - 1)) == 0), + gmx::isPowerOfTwo(c_threadsPerBlock), "Number of threads per block should be a power of two in order for reduction to work."); // This is to prevent the assertion failure for the systems without water diff --git a/src/gromacs/nbnxm/benchmark/bench_system.cpp b/src/gromacs/nbnxm/benchmark/bench_system.cpp index 9ed356933a..999457ae62 100644 --- a/src/gromacs/nbnxm/benchmark/bench_system.cpp +++ b/src/gromacs/nbnxm/benchmark/bench_system.cpp @@ -1,7 +1,7 @@ /* * This file is part of the GROMACS molecular simulation package. * - * Copyright (c) 2019,2020, by the GROMACS development team, led by + * Copyright (c) 2019,2020,2021, 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. @@ -89,7 +89,7 @@ constexpr real c12Oxygen = 2.634129e-06; // A fatal error is generated when this is not the case. static void generateCoordinates(int multiplicationFactor, std::vector* coordinates, matrix box) { - if (multiplicationFactor < 1 || (multiplicationFactor & (multiplicationFactor - 1)) != 0) + if (!gmx::isPowerOfTwo(multiplicationFactor)) { gmx_fatal(FARGS, "The size factor has to be a power of 2"); } diff --git a/src/gromacs/nbnxm/nbnxm_geometry.h b/src/gromacs/nbnxm/nbnxm_geometry.h index 6623885788..68e60974d6 100644 --- a/src/gromacs/nbnxm/nbnxm_geometry.h +++ b/src/gromacs/nbnxm/nbnxm_geometry.h @@ -58,18 +58,12 @@ */ static inline int get_2log(int n) { - // TODO: Replace with gmx::Log2I? - int log2 = 0; - while ((1 << log2) < n) - { - log2++; - } - if ((1 << log2) != n) + if (!gmx::isPowerOfTwo(n)) { gmx_fatal(FARGS, "nbnxn na_c (%d) is not a power of 2", n); } - return log2; + return gmx::log2I(n); } namespace Nbnxm -- 2.22.0