Use constexpr West in src

nliber · Aug 23, 2023 · e6b98b7 · e6b98b7
1 parent 7e299b4
commit e6b98b7
Show file tree

Hide file tree

Showing 9 changed files with 25 additions and 25 deletions.
diff --git a/core/src/HIP/Kokkos_HIP.cpp b/core/src/HIP/Kokkos_HIP.cpp
@@ -60,7 +60,7 @@ void HIP::impl_initialize(InitializationSettings const& settings) {
   if (Impl::HIPTraits::WarpSize < Impl::HIPInternal::m_maxWarpCount) {
     Impl::HIPInternal::m_maxWarpCount = Impl::HIPTraits::WarpSize;
   }
-  int constexpr WordSize              = sizeof(size_type);
+  constexpr int WordSize              = sizeof(size_type);
   Impl::HIPInternal::m_maxSharedWords = hipProp.sharedMemPerBlock / WordSize;
 
   //----------------------------------

diff --git a/core/src/HIP/Kokkos_HIP_Instance.hpp b/core/src/HIP/Kokkos_HIP_Instance.hpp
@@ -31,21 +31,21 @@ namespace Impl {
 struct HIPTraits {
 #if defined(KOKKOS_ARCH_AMD_GFX906) || defined(KOKKOS_ARCH_AMD_GFX908) || \
     defined(KOKKOS_ARCH_AMD_GFX90A) || defined(KOKKOS_ARCH_AMD_GFX942)
-  static int constexpr WarpSize       = 64;
-  static int constexpr WarpIndexMask  = 0x003f; /* hexadecimal for 63 */
-  static int constexpr WarpIndexShift = 6;      /* WarpSize == 1 << WarpShift*/
+  static constexpr int WarpSize       = 64;
+  static constexpr int WarpIndexMask  = 0x003f; /* hexadecimal for 63 */
+  static constexpr int WarpIndexShift = 6;      /* WarpSize == 1 << WarpShift*/
 #elif defined(KOKKOS_ARCH_AMD_GFX1030) || defined(KOKKOS_ARCH_AMD_GFX1100)
-  static int constexpr WarpSize       = 32;
-  static int constexpr WarpIndexMask  = 0x001f; /* hexadecimal for 31 */
-  static int constexpr WarpIndexShift = 5;      /* WarpSize == 1 << WarpShift*/
+  static constexpr int WarpSize       = 32;
+  static constexpr int WarpIndexMask  = 0x001f; /* hexadecimal for 31 */
+  static constexpr int WarpIndexShift = 5;      /* WarpSize == 1 << WarpShift*/
 #endif
-  static int constexpr ConservativeThreadsPerBlock =
+  static constexpr int ConservativeThreadsPerBlock =
       256;  // conservative fallback blocksize in case of spills
-  static int constexpr MaxThreadsPerBlock =
+  static constexpr int MaxThreadsPerBlock =
       1024;  // the maximum we can fit in a block
-  static int constexpr ConstantMemoryUsage        = 0x008000; /* 32k bytes */
-  static int constexpr KernelArgumentLimit        = 0x001000; /*  4k bytes */
-  static int constexpr ConstantMemoryUseThreshold = 0x000200; /* 512 bytes */
+  static constexpr int ConstantMemoryUsage        = 0x008000; /* 32k bytes */
+  static constexpr int KernelArgumentLimit        = 0x001000; /*  4k bytes */
+  static constexpr int ConstantMemoryUseThreshold = 0x000200; /* 512 bytes */
 };
 
 //----------------------------------------------------------------------------

diff --git a/core/src/HIP/Kokkos_HIP_Parallel_Range.hpp b/core/src/HIP/Kokkos_HIP_Parallel_Range.hpp
@@ -150,7 +150,7 @@ class ParallelReduce<CombinedFunctorReducerType, Kokkos::RangePolicy<Traits...>,
   word_size_type* m_scratch_space = nullptr;
   size_type* m_scratch_flags      = nullptr;
 
-  static bool constexpr UseShflReduction = false;
+  static constexpr bool UseShflReduction = false;
 
  private:
   struct ShflReductionTag {};

diff --git a/core/src/HIP/Kokkos_HIP_Parallel_Team.hpp b/core/src/HIP/Kokkos_HIP_Parallel_Team.hpp
@@ -147,7 +147,7 @@ class TeamPolicyInternal<HIP, Properties...>
     // Allow only power-of-two vector_length
     if (!(is_integral_power_of_two(test_vector_length))) {
       int test_pow2           = 1;
-      int constexpr warp_size = HIPTraits::WarpSize;
+      constexpr int warp_size = HIPTraits::WarpSize;
       while (test_pow2 < warp_size) {
         test_pow2 <<= 1;
         if (test_pow2 > test_vector_length) {
@@ -589,7 +589,7 @@ class ParallelReduce<CombinedFunctorReducerType,
  public:
   using size_type = HIP::size_type;
 
-  static int constexpr UseShflReduction =
+  static constexpr int UseShflReduction =
       (ReducerType::static_value_size() != 0);
 
  private:

diff --git a/core/src/HIP/Kokkos_HIP_ReduceScan.hpp b/core/src/HIP/Kokkos_HIP_ReduceScan.hpp
@@ -58,7 +58,7 @@ struct HIPReductionsFunctor<FunctorType, true> {
       FunctorType const& functor, Scalar value, bool const skip,
       Scalar* my_global_team_buffer_element, int const shared_elements,
       Scalar* shared_team_buffer_element) {
-    unsigned int constexpr warp_size = HIPTraits::WarpSize;
+    constexpr unsigned int warp_size = HIPTraits::WarpSize;
     int const warp_id                = (threadIdx.y * blockDim.x) / warp_size;
     Scalar* const my_shared_team_buffer_element =
         shared_team_buffer_element + warp_id % shared_elements;
@@ -105,7 +105,7 @@ struct HIPReductionsFunctor<FunctorType, true> {
     Scalar* shared_team_buffer_elements =
         reinterpret_cast<Scalar*>(shared_data);
     Scalar value                     = shared_team_buffer_elements[threadIdx.y];
-    unsigned int constexpr warp_size = Impl::HIPTraits::WarpSize;
+    constexpr unsigned int warp_size = Impl::HIPTraits::WarpSize;
     int shared_elements              = blockDim.x * blockDim.y / warp_size;
     int global_elements              = block_count;
     __syncthreads();

diff --git a/core/src/HIP/Kokkos_HIP_Shuffle_Reduce.hpp b/core/src/HIP/Kokkos_HIP_Shuffle_Reduce.hpp
@@ -40,7 +40,7 @@ __device__ inline void hip_intra_warp_shuffle_reduction(
   unsigned int shift = 1;
 
   // Reduce over values from threads with different threadIdx.y
-  unsigned int constexpr warp_size = HIPTraits::WarpSize;
+  constexpr unsigned int warp_size = HIPTraits::WarpSize;
   while (blockDim.x * shift < warp_size) {
     ValueType const tmp = shfl_down(result, blockDim.x * shift, warp_size);
     // Only join if upper thread is active (this allows non power of two for
@@ -59,8 +59,8 @@ template <typename ValueType, typename ReducerType>
 __device__ inline void hip_inter_warp_shuffle_reduction(
     ValueType& value, const ReducerType& reducer,
     const int max_active_thread = blockDim.y) {
-  unsigned int constexpr warp_size = HIPTraits::WarpSize;
-  int constexpr step_width         = 8;
+  constexpr unsigned int warp_size = HIPTraits::WarpSize;
+  constexpr int step_width         = 8;
   // Depending on the ValueType __shared__ memory must be aligned up to 8 byte
   // boundaries. The reason not to use ValueType directly is that for types with
   // constructors it could lead to race conditions.
@@ -125,7 +125,7 @@ __device__ inline bool hip_inter_block_shuffle_reduction(
   // block values from global scratch_memory
   bool last_block = false;
   __syncthreads();
-  int constexpr warp_size = HIPTraits::WarpSize;
+  constexpr int warp_size = HIPTraits::WarpSize;
   if (id < warp_size) {
     HIP::size_type count;
 

diff --git a/core/src/HIP/Kokkos_HIP_Vectorization.hpp b/core/src/HIP/Kokkos_HIP_Vectorization.hpp
@@ -85,7 +85,7 @@ struct in_place_shfl_op {
   operator()(Scalar& out, const Scalar& val, int lane_or_delta, int width) const
       noexcept {
     using shuffle_as_t = int;
-    int constexpr N    = sizeof(Scalar) / sizeof(shuffle_as_t);
+    constexpr int N    = sizeof(Scalar) / sizeof(shuffle_as_t);
 
     for (int i = 0; i < N; ++i) {
       reinterpret_cast<shuffle_as_t*>(&out)[i] = self().do_shfl_op(

diff --git a/core/src/Kokkos_Concepts.hpp b/core/src/Kokkos_Concepts.hpp
@@ -117,8 +117,8 @@ template <unsigned int maxT = 0 /* Max threads per block */
 struct LaunchBounds {
   using launch_bounds = LaunchBounds;
   using type          = LaunchBounds<maxT, minB>;
-  static unsigned int constexpr maxTperB{maxT};
-  static unsigned int constexpr minBperSM{minB};
+  static constexpr unsigned int maxTperB{maxT};
+  static constexpr unsigned int minBperSM{minB};
 };
 
 }  // namespace Kokkos

diff --git a/core/src/OpenMP/Kokkos_OpenMP_Instance.hpp b/core/src/OpenMP/Kokkos_OpenMP_Instance.hpp
@@ -48,7 +48,7 @@ class OpenMPInternal;
 inline int g_openmp_hardware_max_threads = 1;
 
 struct OpenMPTraits {
-  static int constexpr MAX_THREAD_COUNT = 512;
+  static constexpr int MAX_THREAD_COUNT = 512;
 };
 
 class OpenMPInternal {