Replace _CCCL_IF_CONSTEXPR (#3775)

NVIDIA · Feb 12, 2025 · 98a3c79 · 98a3c79
1 parent 5857e3d
commit 98a3c79
Show file tree

Hide file tree

Showing 64 changed files with 1,038 additions and 1,063 deletions.
diff --git a/.clang-format b/.clang-format
@@ -87,9 +87,6 @@ ContinuationIndentWidth: 2
 EmptyLineAfterAccessModifier: Never
 EmptyLineBeforeAccessModifier: Always
 FixNamespaceComments: true
-IfMacros: [
-  '_CCCL_IF_CONSTEXPR'
-]
 IndentWrappedFunctionNames: false
 IncludeBlocks:   Regroup
 IncludeCategories:

diff --git a/c2h/include/c2h/catch2_test_helper.h b/c2h/include/c2h/catch2_test_helper.h
@@ -221,7 +221,7 @@ enable_if_t<(N == sizeof...(T))> print_elem(::std::ostream&, const tuple<T...>&)
 template <size_t N, typename... T>
 enable_if_t<(N < sizeof...(T))> print_elem(::std::ostream& os, const tuple<T...>& tup)
 {
-  _CCCL_IF_CONSTEXPR (N != 0)
+  if constexpr (N != 0)
   {
     os << ", ";
   }

diff --git a/cmake/CCCLBuildCompilerTargets.cmake b/cmake/CCCLBuildCompilerTargets.cmake
@@ -162,20 +162,6 @@ function(cccl_build_compiler_targets)
     target_link_libraries(cccl.compiler_interface_cpp${dialect} INTERFACE cccl.compiler_interface)
   endforeach()
 
-  if (CMAKE_CXX_COMPILER_ID STREQUAL "MSVC")
-    # C4127: conditional expression is constant
-    # Disable this MSVC warning for C++11/C++14. In C++17+, we can use
-    # _CCCL_IF_CONSTEXPR to address these warnings.
-    target_compile_options(cccl.compiler_interface_cpp11 INTERFACE
-      $<$<COMPILE_LANGUAGE:CXX>:/wd4127>
-      $<$<COMPILE_LANG_AND_ID:CUDA,NVIDIA>:-Xcompiler=/wd4127>
-    )
-    target_compile_options(cccl.compiler_interface_cpp14 INTERFACE
-      $<$<COMPILE_LANGUAGE:CXX>:/wd4127>
-      $<$<COMPILE_LANG_AND_ID:CUDA,NVIDIA>:-Xcompiler=/wd4127>
-    )
-  endif()
-
   # Some of our unit tests unconditionally throw exceptions, and compilers will
   # detect that the following instructions are unreachable. This is intentional
   # and unavoidable in these cases. This target can be used to silence

diff --git a/cub/cub/agent/agent_merge.cuh b/cub/cub/agent/agent_merge.cuh
@@ -173,7 +173,7 @@ struct agent_t
 
     // if items are provided, merge them
     static constexpr bool have_items = !::cuda::std::is_same_v<item_type, NullType>;
-    _CCCL_IF_CONSTEXPR (have_items)
+    if constexpr (have_items)
     {
       item_type items_loc[items_per_thread];
       merge_sort::gmem_to_reg<threads_per_block, IsFullTile>(

diff --git a/cub/cub/agent/agent_merge_sort.cuh b/cub/cub/agent/agent_merge_sort.cuh
@@ -185,9 +185,9 @@ struct AgentBlockSort
 
     _CCCL_PDL_GRID_DEPENDENCY_SYNC();
 
-    _CCCL_IF_CONSTEXPR (!KEYS_ONLY)
+    if constexpr (!KEYS_ONLY)
     {
-      _CCCL_IF_CONSTEXPR (IS_LAST_TILE)
+      if constexpr (IS_LAST_TILE)
       {
         BlockLoadItems(storage.load_items)
           .Load(items_in + tile_base, items_local, num_remaining, *(items_in + tile_base));
@@ -201,7 +201,7 @@ struct AgentBlockSort
     }
 
     KeyT keys_local[ITEMS_PER_THREAD];
-    _CCCL_IF_CONSTEXPR (IS_LAST_TILE)
+    if constexpr (IS_LAST_TILE)
     {
       BlockLoadKeys(storage.load_keys).Load(keys_in + tile_base, keys_local, num_remaining, *(keys_in + tile_base));
     }
@@ -213,7 +213,7 @@ struct AgentBlockSort
     __syncthreads();
     _CCCL_PDL_TRIGGER_NEXT_LAUNCH();
 
-    _CCCL_IF_CONSTEXPR (IS_LAST_TILE)
+    if constexpr (IS_LAST_TILE)
     {
       BlockMergeSortT(storage.block_merge).Sort(keys_local, items_local, compare_op, num_remaining, keys_local[0]);
     }
@@ -226,7 +226,7 @@ struct AgentBlockSort
 
     if (ping)
     {
-      _CCCL_IF_CONSTEXPR (IS_LAST_TILE)
+      if constexpr (IS_LAST_TILE)
       {
         BlockStoreKeysIt(storage.store_keys_it).Store(keys_out_it + tile_base, keys_local, num_remaining);
       }
@@ -235,11 +235,11 @@ struct AgentBlockSort
         BlockStoreKeysIt(storage.store_keys_it).Store(keys_out_it + tile_base, keys_local);
       }
 
-      _CCCL_IF_CONSTEXPR (!KEYS_ONLY)
+      if constexpr (!KEYS_ONLY)
       {
         __syncthreads();
 
-        _CCCL_IF_CONSTEXPR (IS_LAST_TILE)
+        if constexpr (IS_LAST_TILE)
         {
           BlockStoreItemsIt(storage.store_items_it).Store(items_out_it + tile_base, items_local, num_remaining);
         }
@@ -251,7 +251,7 @@ struct AgentBlockSort
     }
     else
     {
-      _CCCL_IF_CONSTEXPR (IS_LAST_TILE)
+      if constexpr (IS_LAST_TILE)
       {
         BlockStoreKeysRaw(storage.store_keys_raw).Store(keys_out_raw + tile_base, keys_local, num_remaining);
       }
@@ -260,11 +260,11 @@ struct AgentBlockSort
         BlockStoreKeysRaw(storage.store_keys_raw).Store(keys_out_raw + tile_base, keys_local);
       }
 
-      _CCCL_IF_CONSTEXPR (!KEYS_ONLY)
+      if constexpr (!KEYS_ONLY)
       {
         __syncthreads();
 
-        _CCCL_IF_CONSTEXPR (IS_LAST_TILE)
+        if constexpr (IS_LAST_TILE)
         {
           BlockStoreItemsRaw(storage.store_items_raw).Store(items_out_raw + tile_base, items_local, num_remaining);
         }
@@ -391,7 +391,7 @@ template <int BLOCK_THREADS, bool IS_FULL_TILE, int ITEMS_PER_THREAD, class T, c
 _CCCL_DEVICE _CCCL_FORCEINLINE void
 gmem_to_reg(T (&output)[ITEMS_PER_THREAD], It1 input1, It2 input2, int count1, int count2)
 {
-  _CCCL_IF_CONSTEXPR (IS_FULL_TILE)
+  if constexpr (IS_FULL_TILE)
   {
 #pragma unroll
     for (int item = 0; item < ITEMS_PER_THREAD; ++item)
@@ -569,7 +569,7 @@ struct AgentMerge
     //
     ValueT items_local[ITEMS_PER_THREAD];
     (void) items_local; // TODO(bgruber): replace by [[maybe_unused]] in C++17
-    _CCCL_IF_CONSTEXPR (!KEYS_ONLY)
+    if constexpr (!KEYS_ONLY)
     {
       if (ping)
       {
@@ -629,7 +629,7 @@ struct AgentMerge
     // write keys
     if (ping)
     {
-      _CCCL_IF_CONSTEXPR (IS_FULL_TILE)
+      if constexpr (IS_FULL_TILE)
       {
         BlockStoreKeysPing(storage.store_keys_ping).Store(keys_out_ping + tile_base, keys_local);
       }
@@ -640,7 +640,7 @@ struct AgentMerge
     }
     else
     {
-      _CCCL_IF_CONSTEXPR (IS_FULL_TILE)
+      if constexpr (IS_FULL_TILE)
       {
         BlockStoreKeysPong(storage.store_keys_pong).Store(keys_out_pong + tile_base, keys_local);
       }
@@ -651,7 +651,7 @@ struct AgentMerge
     }
 
     // if items are provided, merge them
-    _CCCL_IF_CONSTEXPR (!KEYS_ONLY)
+    if constexpr (!KEYS_ONLY)
     {
       __syncthreads();
 
@@ -673,7 +673,7 @@ struct AgentMerge
       //
       if (ping)
       {
-        _CCCL_IF_CONSTEXPR (IS_FULL_TILE)
+        if constexpr (IS_FULL_TILE)
         {
           BlockStoreItemsPing(storage.store_items_ping).Store(items_out_ping + tile_base, items_local);
         }
@@ -684,7 +684,7 @@ struct AgentMerge
       }
       else
       {
-        _CCCL_IF_CONSTEXPR (IS_FULL_TILE)
+        if constexpr (IS_FULL_TILE)
         {
           BlockStoreItemsPong(storage.store_items_pong).Store(items_out_pong + tile_base, items_local);
         }

diff --git a/cub/cub/agent/agent_sub_warp_merge_sort.cuh b/cub/cub/agent/agent_sub_warp_merge_sort.cuh
@@ -117,7 +117,7 @@ class AgentSubWarpSort
     template <typename T>
     _CCCL_DEVICE bool operator()(T lhs, T rhs) const noexcept
     {
-      _CCCL_IF_CONSTEXPR (IS_DESCENDING)
+      if constexpr (IS_DESCENDING)
       {
         return lhs > rhs;
       }
@@ -132,7 +132,7 @@ class AgentSubWarpSort
     _CCCL_DEVICE bool operator()(__half lhs, __half rhs) const noexcept
     {
       // Need to explicitly cast to float for SM <= 52.
-      _CCCL_IF_CONSTEXPR (IS_DESCENDING)
+      if constexpr (IS_DESCENDING)
       {
         NV_IF_TARGET(NV_PROVIDES_SM_53, (return __hgt(lhs, rhs);), (return __half2float(lhs) > __half2float(rhs);));
       }