Pull out _LIBCUDACXX_UNREACHABLE into its own file (#2399)

* Pull out `_LIBCUDACXX_UNREACHABLE` into its own file Also make it available globally * Drop more uses of `cstdlib`
NVIDIA · Sep 11, 2024 · 1c422f2 · 1c422f2
1 parent e7ade77
commit 1c422f2
Show file tree

Hide file tree

Showing 27 changed files with 108 additions and 106 deletions.
diff --git a/cub/cub/agent/agent_sub_warp_merge_sort.cuh b/cub/cub/agent/agent_sub_warp_merge_sort.cuh
@@ -45,8 +45,6 @@
 
 #include <thrust/system/cuda/detail/core/util.h>
 
-#include <cuda/std/cstdlib>
-
 #include <nv/target>
 
 CUB_NAMESPACE_BEGIN
@@ -122,7 +120,7 @@ class AgentSubWarpSort
       {
         return lhs < rhs;
       }
-      _LIBCUDACXX_UNREACHABLE();
+      _CCCL_UNREACHABLE();
     }
 
 #if defined(__CUDA_FP16_TYPES_EXIST__)
@@ -137,7 +135,7 @@ class AgentSubWarpSort
       {
         NV_IF_TARGET(NV_PROVIDES_SM_53, (return __hlt(lhs, rhs);), (return __half2float(lhs) < __half2float(rhs);));
       }
-      _LIBCUDACXX_UNREACHABLE();
+      _CCCL_UNREACHABLE();
     }
 #endif // __CUDA_FP16_TYPES_EXIST__
   };

diff --git a/cub/test/catch2_test_device_histogram.cu b/cub/test/catch2_test_device_histogram.cu
@@ -305,7 +305,7 @@ void test_even_and_range(LevelT max_level, int max_level_count, OffsetT width, O
       {
         return static_cast<int>((sample - min) * fp_scales[channel]);
       }
-      _LIBCUDACXX_UNREACHABLE();
+      _CCCL_UNREACHABLE();
     };
     auto h_histogram = compute_reference_result<Channels, CounterT>(
       h_samples, sample_to_bin_index, num_levels, width, height, row_pitch);

diff --git a/cudax/include/cuda/experimental/__hierarchy/hierarchy_dimensions.cuh b/cudax/include/cuda/experimental/__hierarchy/hierarchy_dimensions.cuh
@@ -54,7 +54,7 @@ _CCCL_NODISCARD _CCCL_HOST_DEVICE constexpr auto fool_compiler(const dimensions<
   {
     return ex;
   }
-  _LIBCUDACXX_UNREACHABLE();
+  _CCCL_UNREACHABLE();
 }
 
 template <typename QueryLevel, typename Hierarchy>
@@ -96,7 +96,7 @@ struct get_level_helper
     {
       return (*this)(levels...);
     }
-    _LIBCUDACXX_UNREACHABLE();
+    _CCCL_UNREACHABLE();
   }
 };
 } // namespace detail

diff --git a/cudax/include/cuda/experimental/__hierarchy/hierarchy_levels.cuh b/cudax/include/cuda/experimental/__hierarchy/hierarchy_levels.cuh
@@ -292,7 +292,7 @@ template <typename Unit, typename Level>
     return dims_product<typename Level::product_type>(
       extents_impl<SplitLevel, Level>(), extents_impl<Unit, SplitLevel>());
   }
-  _LIBCUDACXX_UNREACHABLE();
+  _CCCL_UNREACHABLE();
 }
 
 template <typename Unit, typename Level>
@@ -309,7 +309,7 @@ template <typename Unit, typename Level>
       dims_product<typename Level::product_type>(index_impl<SplitLevel, Level>(), extents_impl<Unit, SplitLevel>()),
       index_impl<Unit, SplitLevel>());
   }
-  _LIBCUDACXX_UNREACHABLE();
+  _CCCL_UNREACHABLE();
 }
 } // namespace detail
 

diff --git a/libcudacxx/include/cuda/__cccl_config b/libcudacxx/include/cuda/__cccl_config
@@ -21,6 +21,7 @@
 #include <cuda/std/__cccl/ptx_isa.h> // IWYU pragma: export
 #include <cuda/std/__cccl/sequence_access.h> // IWYU pragma: export
 #include <cuda/std/__cccl/system_header.h> // IWYU pragma: export
+#include <cuda/std/__cccl/unreachable.h> // IWYU pragma: export
 #include <cuda/std/__cccl/version.h> // IWYU pragma: export
 #include <cuda/std/__cccl/visibility.h> // IWYU pragma: export
 

diff --git a/libcudacxx/include/cuda/std/__algorithm/copy.h b/libcudacxx/include/cuda/std/__algorithm/copy.h
@@ -28,7 +28,7 @@
 #include <cuda/std/__type_traits/is_trivially_copyable.h>
 #include <cuda/std/__type_traits/remove_const.h>
 #include <cuda/std/cstdint>
-#include <cuda/std/cstdlib>
+#include <cuda/std/cstdlib> // ::memmove
 #include <cuda/std/detail/libcxx/include/cstring>
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD

diff --git a/libcudacxx/include/cuda/std/__cccl/unreachable.h b/libcudacxx/include/cuda/std/__cccl/unreachable.h
@@ -0,0 +1,52 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CCCL_UNREACHABLE_H
+#define __CCCL_UNREACHABLE_H
+
+#include <cuda/std/__cccl/compiler.h>
+#include <cuda/std/__cccl/system_header.h>
+
+#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
+#  pragma GCC system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
+#  pragma clang system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
+#  pragma system_header
+#endif // no system header
+
+#include <cuda/std/__cccl/visibility.h>
+
+#if defined(_CCCL_CUDA_COMPILER_CLANG)
+#  define _CCCL_UNREACHABLE() __builtin_unreachable()
+#elif defined(__CUDA_ARCH__)
+#  if defined(_CCCL_CUDACC_BELOW_11_2)
+#    define _CCCL_UNREACHABLE() __trap()
+#  elif defined(_CCCL_CUDACC_BELOW_11_3)
+#    define _CCCL_UNREACHABLE() __builtin_assume(false)
+#  else
+#    define _CCCL_UNREACHABLE() __builtin_unreachable()
+#  endif // CUDACC above 11.4
+#else // ^^^ __CUDA_ARCH__ ^^^ / vvv !__CUDA_ARCH__ vvv
+#  if defined(_CCCL_COMPILER_MSVC_2017)
+template <class = void>
+_LIBCUDACXX_HIDE_FROM_ABI __declspec(noreturn) void __cccl_unreachable_fallback()
+{
+  __assume(0);
+}
+#    define _CCCL_UNREACHABLE() __cccl_unreachable_fallback()
+#  elif defined(_CCCL_COMPILER_MSVC)
+#    define _CCCL_UNREACHABLE() __assume(0)
+#  else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv
+#    define _CCCL_UNREACHABLE() __builtin_unreachable()
+#  endif // !_CCCL_COMPILER_MSVC
+#endif // !__CUDA_ARCH__
+
+#endif // __CCCL_UNREACHABLE_H
diff --git a/libcudacxx/include/cuda/std/__cuda/barrier.h b/libcudacxx/include/cuda/std/__cuda/barrier.h
@@ -27,7 +27,6 @@
 
 #include <cuda/std/__atomic/api/owned.h>
 #include <cuda/std/__type_traits/void_t.h> // _CUDA_VSTD::void_t
-#include <cuda/std/cstdlib> // _LIBCUDACXX_UNREACHABLE
 
 #if defined(_CCCL_CUDA_COMPILER)
 #  include <cuda/ptx> // cuda::ptx::*
@@ -661,7 +660,7 @@ _CCCL_DEVICE inline async_contract_fulfillment memcpy_async_tx(
         // or from shared to remote cluster dsmem. To copy to remote
         // dsmem, we need to arrive on a cluster-scoped barrier, which
         // is not yet implemented. So we trap in this case as well.
-        _LIBCUDACXX_UNREACHABLE();
+        _CCCL_UNREACHABLE();
       }),
     (__cuda_ptx_memcpy_async_tx_is_not_supported_before_SM_90__();));
 
@@ -782,7 +781,7 @@ struct __memcpy_completion_impl
         // This completion mechanism should not be used with a shared
         // memory barrier. Or at least, we do not currently envision
         // bulk group to be used with shared memory barriers.
-        _LIBCUDACXX_UNREACHABLE();
+        _CCCL_UNREACHABLE();
       case __completion_mechanism::__mbarrier_complete_tx:
 #  if __cccl_ptx_isa >= 800
         // Pre-sm90, the mbarrier_complete_tx completion mechanism is not available.
@@ -798,7 +797,7 @@ struct __memcpy_completion_impl
         return async_contract_fulfillment::none;
       default:
         // Get rid of "control reaches end of non-void function":
-        _LIBCUDACXX_UNREACHABLE();
+        _CCCL_UNREACHABLE();
     }
   }
 
@@ -828,16 +827,16 @@ struct __memcpy_completion_impl
         return async_contract_fulfillment::async;
       case __completion_mechanism::__mbarrier_complete_tx:
         // Non-smem barriers do not have an mbarrier_complete_tx mechanism..
-        _LIBCUDACXX_UNREACHABLE();
+        _CCCL_UNREACHABLE();
       case __completion_mechanism::__async_bulk_group:
         // This completion mechanism is currently not expected to be used with barriers.
-        _LIBCUDACXX_UNREACHABLE();
+        _CCCL_UNREACHABLE();
       case __completion_mechanism::__sync:
         // sync: In this case, we do not need to do anything.
         return async_contract_fulfillment::none;
       default:
         // Get rid of "control reaches end of non-void function":
-        _LIBCUDACXX_UNREACHABLE();
+        _CCCL_UNREACHABLE();
     }
   }
 
@@ -862,7 +861,7 @@ struct __memcpy_completion_impl
         return async_contract_fulfillment::none;
       default:
         // Get rid of "control reaches end of non-void function":
-        _LIBCUDACXX_UNREACHABLE();
+        _CCCL_UNREACHABLE();
     }
   }
 };

diff --git a/libcudacxx/include/cuda/std/__exception/terminate.h b/libcudacxx/include/cuda/std/__exception/terminate.h
@@ -22,11 +22,7 @@
 #  pragma system_header
 #endif // no system header
 
-#if !defined(_CCCL_COMPILER_NVRTC)
-#  include <exception>
-#endif // !_CCCL_COMPILER_NVRTC
-
-#include <cuda/std/cstdlib>
+#include <cuda/std/cstdlib> // ::exit
 
 _CCCL_DIAG_PUSH
 _CCCL_DIAG_SUPPRESS_MSVC(4702) // unreachable code
@@ -35,8 +31,8 @@ _LIBCUDACXX_BEGIN_NAMESPACE_STD_NOVERSION // purposefully not using versioning n
 
 _CCCL_NORETURN _LIBCUDACXX_HIDE_FROM_ABI void __cccl_terminate() noexcept
 {
-  NV_IF_ELSE_TARGET(NV_IS_HOST, (::std::exit(-1);), (__trap();))
-  _LIBCUDACXX_UNREACHABLE();
+  NV_IF_ELSE_TARGET(NV_IS_HOST, (::exit(-1);), (__trap();))
+  _CCCL_UNREACHABLE();
 }
 
 #if 0 // Expose once atomic is universally available
@@ -63,7 +59,7 @@ _LIBCUDACXX_HIDE_FROM_ABI  terminate_handler get_terminate() noexcept
 _CCCL_NORETURN _LIBCUDACXX_HIDE_FROM_ABI void terminate() noexcept
 {
   __cccl_terminate();
-  _LIBCUDACXX_UNREACHABLE();
+  _CCCL_UNREACHABLE();
 }
 
 _LIBCUDACXX_END_NAMESPACE_STD_NOVERSION

diff --git a/libcudacxx/include/cuda/std/__expected/expected.h b/libcudacxx/include/cuda/std/__expected/expected.h
@@ -60,7 +60,6 @@
 #include <cuda/std/__utility/in_place.h>
 #include <cuda/std/__utility/move.h>
 #include <cuda/std/__utility/swap.h>
-#include <cuda/std/cstdlib>
 #include <cuda/std/detail/libcxx/include/__assert>
 #include <cuda/std/initializer_list>
 

diff --git a/libcudacxx/include/cuda/std/__iterator/advance.h b/libcudacxx/include/cuda/std/__iterator/advance.h
@@ -28,7 +28,6 @@
 #include <cuda/std/__iterator/iterator_traits.h>
 #include <cuda/std/__utility/convert_to_integral.h>
 #include <cuda/std/__utility/move.h>
-#include <cuda/std/cstdlib>
 #include <cuda/std/detail/libcxx/include/__assert>
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
@@ -229,7 +228,7 @@ struct __fn
       }
       return __n;
     }
-    _LIBCUDACXX_UNREACHABLE();
+    _CCCL_UNREACHABLE();
   }
 };
 _LIBCUDACXX_END_NAMESPACE_CPO

diff --git a/libcudacxx/include/cuda/std/__iterator/distance.h b/libcudacxx/include/cuda/std/__iterator/distance.h
@@ -92,7 +92,7 @@ struct __fn
     {
       return __last - decay_t<_Ip>(__first);
     }
-    _LIBCUDACXX_UNREACHABLE();
+    _CCCL_UNREACHABLE();
   }
 
   _LIBCUDACXX_TEMPLATE(class _Rp)
@@ -107,7 +107,7 @@ struct __fn
     {
       return operator()(_CUDA_VRANGES::begin(__r), _CUDA_VRANGES::end(__r));
     }
-    _LIBCUDACXX_UNREACHABLE();
+    _CCCL_UNREACHABLE();
   }
 };
 _LIBCUDACXX_END_NAMESPACE_CPO

diff --git a/libcudacxx/include/cuda/std/__iterator/move_iterator.h b/libcudacxx/include/cuda/std/__iterator/move_iterator.h
@@ -41,7 +41,6 @@
 #include <cuda/std/__type_traits/is_reference.h>
 #include <cuda/std/__type_traits/remove_reference.h>
 #include <cuda/std/__utility/move.h>
-#include <cuda/std/cstdlib>
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
@@ -127,7 +126,7 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT move_iterator
     {
       return input_iterator_tag{};
     }
-    _LIBCUDACXX_UNREACHABLE();
+    _CCCL_UNREACHABLE();
   }
 #  endif // !_CCCL_COMPILER_MSVC_2017
 #endif // _CCCL_STD_VER >= 2017

diff --git a/libcudacxx/include/cuda/std/__iterator/reverse_iterator.h b/libcudacxx/include/cuda/std/__iterator/reverse_iterator.h
@@ -271,7 +271,7 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT reverse_iterator
     auto __ytmp = __y.base();
     _CUDA_VRANGES::iter_swap(--__xtmp, --__ytmp);
 #    if defined(_CCCL_COMPILER_ICC)
-    _LIBCUDACXX_UNREACHABLE();
+    _CCCL_UNREACHABLE();
 #    endif // _CCCL_COMPILER_ICC
   }
 #  endif // !_CCCL_COMPILER_MSVC_2017

diff --git a/libcudacxx/include/cuda/std/__memory/allocator.h b/libcudacxx/include/cuda/std/__memory/allocator.h
@@ -32,7 +32,6 @@
 #include <cuda/std/__type_traits/is_volatile.h>
 #include <cuda/std/__utility/forward.h>
 #include <cuda/std/cstddef>
-#include <cuda/std/cstdlib>
 
 #if defined(_CCCL_HAS_CONSTEXPR_ALLOCATION) && !defined(_CCCL_COMPILER_NVRTC)
 #  include <memory>

diff --git a/libcudacxx/include/cuda/std/__ranges/size.h b/libcudacxx/include/cuda/std/__ranges/size.h
@@ -33,7 +33,6 @@
 #include <cuda/std/__utility/auto_cast.h>
 #include <cuda/std/__utility/declval.h>
 #include <cuda/std/cstddef>
-#include <cuda/std/cstdlib>
 
 _LIBCUDACXX_BEGIN_NAMESPACE_RANGES
 
@@ -192,7 +191,7 @@ struct __fn
     {
       return static_cast<_Signed>(_CUDA_VRANGES::size(__t));
     }
-    _LIBCUDACXX_UNREACHABLE();
+    _CCCL_UNREACHABLE();
   }
 };
 _LIBCUDACXX_END_NAMESPACE_CPO

diff --git a/libcudacxx/include/cuda/std/__ranges/subrange.h b/libcudacxx/include/cuda/std/__ranges/subrange.h
@@ -50,7 +50,6 @@
 #include <cuda/std/__type_traits/remove_const.h>
 #include <cuda/std/__type_traits/remove_pointer.h>
 #include <cuda/std/__utility/move.h>
-#include <cuda/std/cstdlib>
 #include <cuda/std/detail/libcxx/include/__assert>
 
 #if _CCCL_STD_VER >= 2017 && !defined(_CCCL_COMPILER_MSVC_2017)
@@ -345,7 +344,7 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT subrange : public view_interface<subrange<_I
       return _CUDA_VSTD::__to_unsigned_like(__end_ - __begin_);
     }
 #  if defined(_CCCL_CUDACC_BELOW_11_3)
-    _LIBCUDACXX_UNREACHABLE();
+    _CCCL_UNREACHABLE();
 #  endif // _CCCL_CUDACC_BELOW_11_3
   }
 
@@ -444,7 +443,7 @@ _LIBCUDACXX_HIDE_FROM_ABI constexpr auto get(const subrange<_Iter, _Sent, _Kind>
   {
     return __subrange.end();
   }
-  _LIBCUDACXX_UNREACHABLE();
+  _CCCL_UNREACHABLE();
 }
 
 #  if _CCCL_STD_VER >= 2020
@@ -468,7 +467,7 @@ _LIBCUDACXX_HIDE_FROM_ABI constexpr auto get(subrange<_Iter, _Sent, _Kind>&& __s
   {
     return __subrange.end();
   }
-  _LIBCUDACXX_UNREACHABLE();
+  _CCCL_UNREACHABLE();
 }
 
 template <class _Ip, class _Sp, subrange_kind _Kp>

diff --git a/libcudacxx/include/cuda/std/__ranges/unwrap_end.h b/libcudacxx/include/cuda/std/__ranges/unwrap_end.h
@@ -43,7 +43,7 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr iterator_t<_Range> __unwrap_
     _CUDA_VRANGES::advance(__ret, _CUDA_VRANGES::end(__range));
     return __ret;
   }
-  _LIBCUDACXX_UNREACHABLE();
+  _CCCL_UNREACHABLE();
 }
 
 #endif // _CCCL_STD_VER >= 2017 && !_CCCL_COMPILER_MSVC_2017

diff --git a/libcudacxx/include/cuda/std/__utility/unreachable.h b/libcudacxx/include/cuda/std/__utility/unreachable.h
@@ -20,20 +20,18 @@
 #  pragma system_header
 #endif // no system header
 
-#include <cuda/std/cstdlib>
-
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
 _CCCL_NORETURN _LIBCUDACXX_HIDE_FROM_ABI void __libcpp_unreachable()
 {
-  _LIBCUDACXX_UNREACHABLE();
+  _CCCL_UNREACHABLE();
 }
 
 #if _CCCL_STD_VER > 2020
 
 [[noreturn]] _LIBCUDACXX_HIDE_FROM_ABI void unreachable()
 {
-  _LIBCUDACXX_UNREACHABLE();
+  _CCCL_UNREACHABLE();
 }
 
 #endif // _CCCL_STD_VER > 2020