Add gmx::isPowerOfTwo function

A follow-up from a discussion in !1080.

diff --git a/api/legacy/include/gromacs/math/functions.h b/api/legacy/include/gromacs/math/functions.h
index 271e1cf76035529d27aa07b4fb5313bbb2c62303..af557fb6adde0cd37262f2b659a1ca92de888c82 100644 (file)
--- 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<int>(rint(x));
 }
 
+//! \brief Check whether \p v is an integer power of 2.
+template<typename T, typename = std::enable_if_t<std::is_integral<T>::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 a81791765c73f44c20ec3c6b5971e368ae586f29..35111b52ec18fdd7e7f2a6aed828f28a7961f857 100644 (file)
--- 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 88e529b553e1a5601e8bfd33bb946b20f3e744fe..03ca35bfeee4c1ba085e00a1ac4a308344c584da 100644 (file)
--- 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<unsigned int factor>
 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 b444ee78412448af973b9f75c279a9002f4314c0..939a08da2eb5f5c8fee3575e1c6161cc7cd0466b 100644 (file)
--- 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<typename T>
+class FunctionTestIntegerTypes : public ::testing::Test
+{
+};
+
+typedef ::testing::Types<char, unsigned char, int, unsigned int, long, unsigned long> IntegerTypes;
+TYPED_TEST_CASE(FunctionTestIntegerTypes, IntegerTypes);
+
+TYPED_TEST(FunctionTestIntegerTypes, IsPowerOfTwo)
+{
+    if (std::is_signed_v<TypeParam>)
+    {
+        EXPECT_EQ(false, gmx::isPowerOfTwo<TypeParam>(std::numeric_limits<TypeParam>::min()));
+        EXPECT_EQ(false, gmx::isPowerOfTwo<TypeParam>(-16));
+        EXPECT_EQ(false, gmx::isPowerOfTwo<TypeParam>(-3));
+        EXPECT_EQ(false, gmx::isPowerOfTwo<TypeParam>(-2));
+        EXPECT_EQ(false, gmx::isPowerOfTwo<TypeParam>(-1));
+    }
+    EXPECT_EQ(false, gmx::isPowerOfTwo<TypeParam>(0));
+    EXPECT_EQ(true, gmx::isPowerOfTwo<TypeParam>(1));
+    EXPECT_EQ(true, gmx::isPowerOfTwo<TypeParam>(2));
+    EXPECT_EQ(false, gmx::isPowerOfTwo<TypeParam>(3));
+    EXPECT_EQ(true, gmx::isPowerOfTwo<TypeParam>(4));
+    EXPECT_EQ(false, gmx::isPowerOfTwo<TypeParam>(5));
+    EXPECT_EQ(false, gmx::isPowerOfTwo<TypeParam>(6));
+    EXPECT_EQ(false, gmx::isPowerOfTwo<TypeParam>(24));
+    EXPECT_EQ(false, gmx::isPowerOfTwo<TypeParam>(63));
+    EXPECT_EQ(true, gmx::isPowerOfTwo<TypeParam>(64));
+    EXPECT_EQ(false, gmx::isPowerOfTwo<TypeParam>(66));
+    // Max for any type is always 2^x - 1
+    EXPECT_EQ(false, gmx::isPowerOfTwo<TypeParam>(std::numeric_limits<TypeParam>::max()));
+}
+
 } // namespace

diff --git a/src/gromacs/mdlib/lincs_gpu.cu b/src/gromacs/mdlib/lincs_gpu.cu
index ac879bcfaed5e7ed48cf68010ee07b5cad5854ed..0967c20781b0470c595c3274d59f7eca877ad852 100644 (file)
--- 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 eb4f7b066b77dfa449678cf59aff7e8460eca1a9..eb4b4957724c5e56733f79358f1b4515a208aeba 100644 (file)
--- 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 9ed356933a67b310f30718eca5007b774c57f85d..999457ae62a39de37d50f946dbce571fe481c346 100644 (file)
--- 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<gmx::RVec>* 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 66238857881427db03fd644c3149591c4a0aef43..68e60974d6f133a964af97afe70c80b7b8a3787c 100644 (file)
--- 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