diff --git a/cudax/include/cuda/experimental/__async/meta.cuh b/cudax/include/cuda/experimental/__async/meta.cuh
index 1de340e51c..3ad9211e95 100644
--- a/cudax/include/cuda/experimental/__async/meta.cuh
+++ b/cudax/include/cuda/experimental/__async/meta.cuh
@@ -534,13 +534,13 @@ using __m_at = _Ts...[__v<_Np>];
 template <size_t _Np, class... _Ts>
 using __m_at_c = _Ts...[_Np];
 
-#elif __has_builtin(__type_pack_element)
+#elif defined(_CCCL_BUILTIN_TYPE_PACK_ELEMENT)
 
 template <bool>
 struct __m_at_
 {
   template <class _Np, class... _Ts>
-  using __f = __type_pack_element<__v<_Np>, _Ts...>;
+  using __f = _CCCL_BUILTIN_TYPE_PACK_ELEMENT(__v<_Np>, _Ts...);
 };
 
 template <class _Np, class... _Ts>
diff --git a/cudax/include/cuda/experimental/__async/type_traits.cuh b/cudax/include/cuda/experimental/__async/type_traits.cuh
index 2fcdb78fde..be15ad87db 100644
--- a/cudax/include/cuda/experimental/__async/type_traits.cuh
+++ b/cudax/include/cuda/experimental/__async/type_traits.cuh
@@ -30,26 +30,13 @@
 
 namespace cuda::experimental::__async
 {
-#if __has_builtin(__remove_reference)
 
 template <class _Ty>
-using __remove_ref_t = __remove_reference(_Ty);
-
-#elif __has_builtin(__remove_reference_t)
-
-template <class _Ty>
-using __remove_ref_t = __remove_reference_t(_Ty);
-
-#else
-
-template <class _Ty>
-using __remove_ref_t = _CUDA_VSTD::remove_reference_t<_Ty>;
-
-#endif
+using __remove_ref_t = _CUDA_VSTD::__libcpp_remove_reference_t<_Ty>;
 
 //////////////////////////////////////////////////////////////////////////////////////////////////
 // __decay_t: An efficient implementation for ::std::decay
-#if __has_builtin(__decay)
+#if defined(_CCCL_BUILTIN_DECAY)
 
 template <class _Ty>
 using __decay_t = __decay(_Ty);
@@ -59,7 +46,7 @@ using __decay_t = __decay(_Ty);
 //   template <class _Ty>
 //   using __decay_t = _CUDA_VSTD::decay_t<_Ty>;
 
-#else
+#else // ^^^ _CCCL_BUILTIN_DECAY ^^^ / vvv !_CCCL_BUILTIN_DECAY vvv
 
 struct __decay_object
 {
@@ -141,7 +128,7 @@ inline __decay_void __mdecay<void const>;
 template <class _Ty>
 using __decay_t = typename decltype(__mdecay<_Ty>)::template __f<_Ty>;
 
-#endif
+#endif // _CCCL_BUILTIN_DECAY
 
 //////////////////////////////////////////////////////////////////////////////////////////////////
 // __copy_cvref_t: For copying cvref from one type to another
diff --git a/libcudacxx/include/cuda/std/__algorithm/copy.h b/libcudacxx/include/cuda/std/__algorithm/copy.h
index f3c4e2874a..bbd7d60e98 100644
--- a/libcudacxx/include/cuda/std/__algorithm/copy.h
+++ b/libcudacxx/include/cuda/std/__algorithm/copy.h
@@ -91,13 +91,13 @@ template <class _Tp, class _Up>
 _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 bool __constexpr_tail_overlap(_Tp* __first, _Up* __needle, _Tp* __last)
 {
   _LIBCUDACXX_UNUSED_VAR(__last);
-#if __has_builtin(__builtin_constant_p) || defined(_CCCL_COMPILER_GCC)
+#if defined(_CCCL_BUILTIN_CONSTANT_P)
   NV_IF_ELSE_TARGET(NV_IS_HOST,
-                    (return __builtin_constant_p(__first < __needle) && __first < __needle;),
+                    (return _CCCL_BUILTIN_CONSTANT_P(__first < __needle) && __first < __needle;),
                     (return __constexpr_tail_overlap_fallback(__first, __needle, __last);))
-#else
+#else // ^^^ _CCCL_BUILTIN_CONSTANT_P ^^^ / vvv !_CCCL_BUILTIN_CONSTANT_P vvv
   return __constexpr_tail_overlap_fallback(__first, __needle, __last);
-#endif
+#endif // !_CCCL_BUILTIN_CONSTANT_P
 }
 
 template <class _AlgPolicy,
diff --git a/libcudacxx/include/cuda/std/__cccl/builtin.h b/libcudacxx/include/cuda/std/__cccl/builtin.h
index dafdc16f83..0a6c5534f4 100644
--- a/libcudacxx/include/cuda/std/__cccl/builtin.h
+++ b/libcudacxx/include/cuda/std/__cccl/builtin.h
@@ -46,16 +46,42 @@
 // MSVC needs manual handling, has no real way of checking builtins so all is manual
 // GCC  needs manual handling, before gcc-10 as that finally supports __has_builtin
 
-#if __check_builtin(array_rank)
+// We enable newer clang builtins through __has_builtin, as that avoids some compiler issues with older gcc / nvcc
+
+// NVCC has issues with function pointers
+#if __has_builtin(__add_lvalue_reference) && defined(_CCCL_CUDA_COMPILER_CLANG)
+#  define _CCCL_BUILTIN_ADD_LVALUE_REFERENCE(...) __add_lvalue_reference(__VA_ARGS__)
+#endif // __has_builtin(__add_lvalue_reference)
+
+// NVCC has issues with function pointers
+#if __has_builtin(__add_pointer) && defined(_CCCL_CUDA_COMPILER_CLANG)
+#  define _CCCL_BUILTIN_ADD_POINTER(...) __add_pointer(__VA_ARGS__)
+#endif // __has_builtin(__add_pointer)
+
+// NVCC has issues with function pointers
+#if __has_builtin(__add_rvalue_reference) && defined(_CCCL_CUDA_COMPILER_CLANG)
+#  define _CCCL_BUILTIN_ADD_RVALUE_REFERENCE(...) __add_rvalue_reference(__VA_ARGS__)
+#endif // __has_builtin(__add_rvalue_reference)
+
+// TODO: Enable using the builtin __array_rank when https://llvm.org/PR57133 is resolved
+#if 0 // __check_builtin(array_rank)
 #  define _CCCL_BUILTIN_ARRAY_RANK(...) __array_rank(__VA_ARGS__)
 #endif // __check_builtin(array_rank)
 
+#if __has_builtin(__array_extent)
+#  define _CCCL_BUILTIN_ARRAY_EXTENT(...) __array_extent(__VA_ARGS__)
+#endif // __has_builtin(array_extent)
+
 // nvhpc has a bug where it supports __builtin_addressof but does not mark it via __check_builtin
 #if __check_builtin(builtin_addressof) || (defined(_CCCL_COMPILER_GCC) && _CCCL_GCC_VERSION >= 70000) \
   || defined(_CCCL_COMPILER_MSVC) || defined(_CCCL_COMPILER_NVHPC)
 #  define _CCCL_BUILTIN_ADDRESSOF(...) __builtin_addressof(__VA_ARGS__)
 #endif // __check_builtin(builtin_addressof)
 
+#if __check_builtin(builtin_assume) && !defined(_CCCL_CUDACC_BELOW_11_2)
+#  define _CCCL_BUILTIN_ASSUME(...) __builtin_assume(__VA_ARGS__)
+#endif // __check_builtin(builtin_assume) && nvcc >= 11.2
+
 // MSVC supports __builtin_bit_cast from 19.25 on
 // clang-9 supports __builtin_bit_cast but it is not a constant expression
 #if (__check_builtin(builtin_bit_cast) || (defined(_CCCL_COMPILER_MSVC) && _CCCL_MSVC_VERSION > 1925)) \
@@ -63,6 +89,10 @@
 #  define _CCCL_BUILTIN_BIT_CAST(...) __builtin_bit_cast(__VA_ARGS__)
 #endif // __check_builtin(builtin_bit_cast)
 
+#if __check_builtin(builtin_contant_p) || defined(_CCCL_COMPILER_GCC)
+#  define _CCCL_BUILTIN_CONSTANT_P(...) __builtin_constant_p(__VA_ARGS__)
+#endif // __check_builtin(builtin_contant_p)
+
 #if __check_builtin(builtin_is_constant_evaluated) || (defined(_CCCL_COMPILER_GCC) && _CCCL_GCC_VERSION >= 90000) \
   || (defined(_CCCL_COMPILER_MSVC) && _CCCL_MSVC_VERSION > 1924 && !defined(_CCCL_CUDACC_BELOW_11_3))
 #  define _CCCL_BUILTIN_IS_CONSTANT_EVALUATED(...) __builtin_is_constant_evaluated(__VA_ARGS__)
@@ -74,84 +104,95 @@
 #  undef _CCCL_BUILTIN_IS_CONSTANT_EVALUATED
 #endif // _CCCL_STD_VER < 2014 && _CCCL_CUDA_COMPILER_NVCC
 
-#if __check_builtin(builtin_launder) || (defined(_CCCL_COMPILER_GCC) && _CCCL_GCC_VERSION >= 70000)
+#if (__check_builtin(builtin_launder) || (defined(_CCCL_COMPILER_GCC) && _CCCL_GCC_VERSION >= 70000)) \
+  && (!defined(_CCCL_COMPILER_CLANG) || _CCCL_CLANG_VERSION >= 100000 || !defined(_CCCL_CUDACC_BELOW_11_3))
 #  define _CCCL_BUILTIN_LAUNDER(...) __builtin_launder(__VA_ARGS__)
-#endif // __check_builtin(builtin_launder)
+#endif // __check_builtin(builtin_launder) && gcc >= 7
 
-// Disabled due to libstdc++ conflict
-#if 0 // __check_builtin(decay)
+#if __check_builtin(__builtin_operator_new) && __check_builtin(__builtin_operator_delete) \
+  && defined(_CCCL_CUDA_COMPILER_CLANG)
+#  define _CCCL_BUILTIN_OPERATOR_DELETE(...) __builtin_operator_delete(__VA_ARGS__)
+#  define _CCCL_BUILTIN_OPERATOR_NEW(...)    __builtin_operator_new(__VA_ARGS__)
+#endif // __check_builtin(__builtin_operator_new) && __check_builtin(__builtin_operator_delete)
+
+#if __has_builtin(__decay) && defined(_CCCL_CUDA_COMPILER_CLANG)
 #  define _CCCL_BUILTIN_DECAY(...) __decay(__VA_ARGS__)
-#endif // __check_builtin(decay)
+#endif // __has_builtin(__decay) && clang-cuda
 
 #if __check_builtin(has_nothrow_assign) || (defined(_CCCL_COMPILER_GCC) && _CCCL_GCC_VERSION >= 40300) \
   || defined(_CCCL_COMPILER_MSVC) || defined(_CCCL_COMPILER_NVRTC)
 #  define _CCCL_BUILTIN_HAS_NOTHROW_ASSIGN(...) __has_nothrow_assign(__VA_ARGS__)
-#endif // __check_builtin(has_nothrow_assign)
+#endif // __check_builtin(has_nothrow_assign) && gcc >= 4.3
 
 #if __check_builtin(has_nothrow_constructor) || (defined(_CCCL_COMPILER_GCC) && _CCCL_GCC_VERSION >= 40300) \
   || defined(_CCCL_COMPILER_MSVC) || defined(_CCCL_COMPILER_NVRTC)
 #  define _CCCL_BUILTIN_HAS_NOTHROW_CONSTRUCTOR(...) __has_nothrow_constructor(__VA_ARGS__)
-#endif // __check_builtin(has_nothrow_constructor)
+#endif // __check_builtin(has_nothrow_constructor) && gcc >= 4.3
 
 #if __check_builtin(has_nothrow_copy) || (defined(_CCCL_COMPILER_GCC) && _CCCL_GCC_VERSION >= 40300) \
   || defined(_CCCL_COMPILER_MSVC) || defined(_CCCL_COMPILER_NVRTC)
 #  define _CCCL_BUILTIN_HAS_NOTHROW_COPY(...) __has_nothrow_copy(__VA_ARGS__)
-#endif // __check_builtin(has_nothrow_copy)
+#endif // __check_builtin(has_nothrow_copy) && gcc >= 4.3
 
 #if __check_builtin(has_trivial_constructor) || (defined(_CCCL_COMPILER_GCC) && _CCCL_GCC_VERSION >= 40300) \
   || defined(_CCCL_COMPILER_MSVC) || defined(_CCCL_COMPILER_NVRTC)
 #  define _CCCL_BUILTIN_HAS_TRIVIAL_CONSTRUCTOR(...) __has_trivial_constructor(__VA_ARGS__)
-#endif // __check_builtin(has_trivial_constructor)
+#endif // __check_builtin(has_trivial_constructor) && gcc >= 4.3
 
 #if __check_builtin(has_trivial_destructor) || (defined(_CCCL_COMPILER_GCC) && _CCCL_GCC_VERSION >= 40300) \
   || defined(_CCCL_COMPILER_MSVC) || defined(_CCCL_COMPILER_NVRTC)
 #  define _CCCL_BUILTIN_HAS_TRIVIAL_DESTRUCTOR(...) __has_trivial_destructor(__VA_ARGS__)
-#endif // __check_builtin(has_trivial_destructor)
+#endif // __check_builtin(has_trivial_destructor) && gcc >= 4.3
 
 #if __check_builtin(has_unique_object_representations) || (defined(_CCCL_COMPILER_GCC) && _CCCL_GCC_VERSION >= 70000)
 #  define _CCCL_BUILTIN_HAS_UNIQUE_OBJECT_REPRESENTATIONS(...) __has_unique_object_representations(__VA_ARGS__)
-#endif // __check_builtin(has_unique_object_representations)
+#endif // __check_builtin(has_unique_object_representations) && gcc >= 7.0
 
 #if __check_builtin(has_virtual_destructor) || (defined(_CCCL_COMPILER_GCC) && _CCCL_GCC_VERSION >= 40300) \
   || defined(_CCCL_COMPILER_MSVC) || defined(_CCCL_COMPILER_NVRTC)
 #  define _CCCL_BUILTIN_HAS_VIRTUAL_DESTRUCTOR(...) __has_virtual_destructor(__VA_ARGS__)
-#endif // __check_builtin(has_virtual_destructor)
+#endif // __check_builtin(has_virtual_destructor) && gcc >= 4.3
+
+#if __has_builtin(__integer_pack)
+#  define _CCCL_BUILTIN_INTEGER_PACK(...) __integer_pack(__VA_ARGS__)
+#endif // __has_builtin(__integer_pack)
 
 #if __check_builtin(is_aggregate) || (defined(_CCCL_COMPILER_GCC) && _CCCL_GCC_VERSION >= 70000) \
   || (defined(_CCCL_COMPILER_MSVC) && _CCCL_MSVC_VERSION > 1914) || defined(_CCCL_COMPILER_NVRTC)
 #  define _CCCL_BUILTIN_IS_AGGREGATE(...) __is_aggregate(__VA_ARGS__)
-#endif // __check_builtin(is_aggregate)
+#endif // __check_builtin(is_aggregate) && gcc >= 7.0
 
-#if __check_builtin(is_array)
+#if __check_builtin(is_array) && (!defined(_CCCL_COMPILER_CLANG) || _CCCL_CLANG_VERSION >= 190000)
 #  define _CCCL_BUILTIN_IS_ARRAY(...) __is_array(__VA_ARGS__)
-#endif // __check_builtin(is_array)
+#endif // __check_builtin(is_array) && clang >= 19
 
-// TODO: Clang incorrectly reports that __is_array is true for T[0].
-//       Re-enable the branch once https://llvm.org/PR54705 is fixed.
-#ifndef _LIBCUDACXX_USE_IS_ARRAY_FALLBACK
-#  if defined(_CCCL_COMPILER_CLANG)
-#    define _LIBCUDACXX_USE_IS_ARRAY_FALLBACK
-#  endif // _CCCL_COMPILER_CLANG
-#endif // !_LIBCUDACXX_USE_IS_ARRAY_FALLBACK
-
-#if __check_builtin(is_assignable) || defined(_CCCL_COMPILER_MSVC)
+#if (__check_builtin(is_assignable) || defined(_CCCL_COMPILER_MSVC)) \
+  && (!defined(_CCCL_COMPILER_GCC) || _CCCL_GCC_VERSION >= 90000)
 #  define _CCCL_BUILTIN_IS_ASSIGNABLE(...) __is_assignable(__VA_ARGS__)
-#endif // __check_builtin(is_assignable)
+#endif // __check_builtin(is_assignable) && gcc >= 9.0
 
 #if __check_builtin(is_base_of) || (defined(_CCCL_COMPILER_GCC) && _CCCL_GCC_VERSION >= 40300) \
   || defined(_CCCL_COMPILER_MSVC) || defined(_CCCL_COMPILER_NVRTC)
 #  define _CCCL_BUILTIN_IS_BASE_OF(...) __is_base_of(__VA_ARGS__)
-#endif // __check_builtin(is_base_of)
+#endif // __check_builtin(is_base_of) && gcc >= 4.3
 
 #if __check_builtin(is_class) || (defined(_CCCL_COMPILER_GCC) && _CCCL_GCC_VERSION >= 40300) \
   || defined(_CCCL_COMPILER_MSVC) || defined(_CCCL_COMPILER_NVRTC)
 #  define _CCCL_BUILTIN_IS_CLASS(...) __is_class(__VA_ARGS__)
-#endif // __check_builtin(is_class)
+#endif // __check_builtin(is_class) && gcc >= 4.3
+
+#if __has_builtin(__is_compound)
+#  define _CCCL_BUILTIN_IS_COMPOUND(...) __is_compound(__VA_ARGS__)
+#endif // __has_builtin(__is_compound)
+
+#if __has_builtin(__is_const)
+#  define _CCCL_BUILTIN_IS_CONST(...) __is_const(__VA_ARGS__)
+#endif // __has_builtin(__is_const)
 
 #if __check_builtin(is_constructible) || (defined(_CCCL_COMPILER_GCC) && _CCCL_GCC_VERSION >= 80000) \
   || defined(_CCCL_COMPILER_MSVC) || defined(_CCCL_COMPILER_NVRTC)
 #  define _CCCL_BUILTIN_IS_CONSTRUCTIBLE(...) __is_constructible(__VA_ARGS__)
-#endif // __check_builtin(is_constructible)
+#endif // __check_builtin(is_constructible) && gcc >= 8.0
 
 #if __check_builtin(is_convertible_to) || defined(_CCCL_COMPILER_MSVC) || defined(_CCCL_COMPILER_NVRTC)
 #  define _CCCL_BUILTIN_IS_CONVERTIBLE_TO(...) __is_convertible_to(__VA_ARGS__)
@@ -164,36 +205,50 @@
 #if __check_builtin(is_empty) || (defined(_CCCL_COMPILER_GCC) && _CCCL_GCC_VERSION >= 40300) \
   || defined(_CCCL_COMPILER_MSVC) || defined(_CCCL_COMPILER_NVRTC)
 #  define _CCCL_BUILTIN_IS_EMPTY(...) __is_empty(__VA_ARGS__)
-#endif // __check_builtin(is_empty)
+#endif // __check_builtin(is_empty) && gcc >= 4.3
 
 #if __check_builtin(is_enum) || (defined(_CCCL_COMPILER_GCC) && _CCCL_GCC_VERSION >= 40300) \
   || defined(_CCCL_COMPILER_MSVC) || defined(_CCCL_COMPILER_NVRTC)
 #  define _CCCL_BUILTIN_IS_ENUM(...) __is_enum(__VA_ARGS__)
-#endif // __check_builtin(is_enum)
+#endif // __check_builtin(is_enum) && gcc >= 4.3
 
 #if __check_builtin(is_final) || (defined(_CCCL_COMPILER_GCC) && _CCCL_GCC_VERSION >= 40700) \
   || defined(_CCCL_COMPILER_MSVC) || defined(_CCCL_COMPILER_NVRTC)
 #  define _CCCL_BUILTIN_IS_FINAL(...) __is_final(__VA_ARGS__)
-#endif // __check_builtin(is_final)
+#endif // __check_builtin(is_final) && gcc >= 4.7
 
 #if __check_builtin(is_function) && !defined(_CCCL_CUDA_COMPILER_NVCC)
 #  define _CCCL_BUILTIN_IS_FUNCTION(...) __is_function(__VA_ARGS__)
 #endif // __check_builtin(is_function)
 
+#if __check_builtin(is_fundamental) && (!defined(_CCCL_COMPILER_CLANG) || _CCCL_CLANG_VERSION >= 100000)
+#  define _CCCL_BUILTIN_IS_FUNDAMENTAL(...) __is_fundamental(__VA_ARGS__)
+#endif // __check_builtin(is_fundamental) && clang >= 10
+
+#if __has_builtin(__is_integral)
+#  define _CCCL_BUILTIN_IS_INTEGRAL(...) __is_integral(__VA_ARGS__)
+#endif // __has_builtin(__is_integral) && nvcc >= 11.3
+
 #if __check_builtin(is_literal_type) || (defined(_CCCL_COMPILER_GCC) && _CCCL_GCC_VERSION >= 40600) \
   || defined(_CCCL_COMPILER_MSVC) || defined(_CCCL_COMPILER_NVRTC)
 #  define _CCCL_BUILTIN_IS_LITERAL(...) __is_literal_type(__VA_ARGS__)
-#endif // __check_builtin(is_literal_type)
+#endif // __check_builtin(is_literal_type) && gcc >= 4.6
 
-#if __check_builtin(is_lvalue_reference)
+#if __check_builtin(is_lvalue_reference) && !defined(_CCCL_CUDACC_BELOW_11_3)
 #  define _CCCL_BUILTIN_IS_LVALUE_REFERENCE(...) __is_lvalue_reference(__VA_ARGS__)
-#endif // __check_builtin(is_lvalue_reference)
+#endif // __check_builtin(is_lvalue_reference) && nvcc >= 11.3
 
-#ifndef _LIBCUDACXX_USE_IS_LVALUE_REFERENCE_FALLBACK
-#  if defined(_CCCL_CUDACC_BELOW_11_3)
-#    define _LIBCUDACXX_USE_IS_LVALUE_REFERENCE_FALLBACK
-#  endif // nvcc < 11.3
-#endif // !_LIBCUDACXX_USE_IS_LVALUE_REFERENCE_FALLBACK
+#if __has_builtin(__is_member_function_pointer)
+#  define _CCCL_BUILTIN_IS_MEMBER_FUNCTION_POINTER(...) __is_member_function_pointer(__VA_ARGS__)
+#endif // __has_builtin(__is_member_function_pointer)
+
+#if __has_builtin(__is_member_object_pointer)
+#  define _CCCL_BUILTIN_IS_MEMBER_OBJECT_POINTER(...) __is_member_object_pointer(__VA_ARGS__)
+#endif // __has_builtin(__is_member_object_pointer)
+
+#if __has_builtin(__is_member_pointer)
+#  define _CCCL_BUILTIN_IS_MEMBER_POINTER(...) __is_member_pointer(__VA_ARGS__)
+#endif // __has_builtin(__is_member_pointer)
 
 #if __check_builtin(is_nothrow_assignable) || defined(_CCCL_COMPILER_MSVC) || defined(_CCCL_COMPILER_NVRTC)
 #  define _CCCL_BUILTIN_IS_NOTHROW_ASSIGNABLE(...) __is_nothrow_assignable(__VA_ARGS__)
@@ -207,82 +262,76 @@
 #  define _CCCL_BUILTIN_IS_NOTHROW_DESTRUCTIBLE(...) __is_nothrow_destructible(__VA_ARGS__)
 #endif // __check_builtin(is_nothrow_destructible)
 
-#if __check_builtin(is_object)
+#if __check_builtin(is_object) && !defined(_CCCL_CUDACC_BELOW_11_3)
 #  define _CCCL_BUILTIN_IS_OBJECT(...) __is_object(__VA_ARGS__)
-#endif // __check_builtin(is_object)
-
-#ifndef _LIBCUDACXX_USE_IS_OBJECT_FALLBACK
-#  if defined(_CCCL_CUDACC_BELOW_11_3)
-#    define _LIBCUDACXX_USE_IS_OBJECT_FALLBACK
-#  endif // _CCCL_CUDACC_BELOW_11_3
-#endif // !_LIBCUDACXX_USE_IS_OBJECT_FALLBACK
+#endif // __check_builtin(is_object) && nvcc >= 11.3
 
 #if __check_builtin(is_pod) || (defined(_CCCL_COMPILER_GCC) && _CCCL_GCC_VERSION >= 40300) \
   || defined(_CCCL_COMPILER_MSVC) || defined(_CCCL_COMPILER_NVRTC)
 #  define _CCCL_BUILTIN_IS_POD(...) __is_pod(__VA_ARGS__)
-#endif // __check_builtin(is_pod)
+#endif // __check_builtin(is_pod) && gcc >= 4.3
 
-// libstdc++ defines this as a function, breaking functionality
-#if 0 // __check_builtin(is_pointer)
+// Disabled due to libstdc++ conflict
+#if 0 // __has_builtin(__is_pointer)
 #  define _CCCL_BUILTIN_IS_POINTER(...) __is_pointer(__VA_ARGS__)
-#endif // __check_builtin(is_pointer)
+#endif // __has_builtin(__is_pointer)
 
 #if __check_builtin(is_polymorphic) || (defined(_CCCL_COMPILER_GCC) && _CCCL_GCC_VERSION >= 40300) \
   || defined(_CCCL_COMPILER_MSVC) || defined(_CCCL_COMPILER_NVRTC)
 #  define _CCCL_BUILTIN_IS_POLYMORPHIC(...) __is_polymorphic(__VA_ARGS__)
-#endif // __check_builtin(is_polymorphic)
+#endif // __check_builtin(is_polymorphic) && gcc >= 4.3
 
-#if __check_builtin(is_reference)
+#if __has_builtin(__is_reference)
 #  define _CCCL_BUILTIN_IS_REFERENCE(...) __is_reference(__VA_ARGS__)
-#endif // __check_builtin(is_reference)
+#endif // __has_builtin(__is_reference)
 
 // Disabled due to libstdc++ conflict
-#if 0 // __check_builtin(is_referenceable)
+#if 0 // __has_builtin(__is_referenceable)
 #  define _CCCL_BUILTIN_IS_REFERENCEABLE(...) __is_referenceable(__VA_ARGS__)
-#endif // __check_builtin(is_referenceable)
+#endif // __has_builtin(__is_referenceable)
 
-#if __check_builtin(is_rvalue_reference)
+#if __has_builtin(__is_rvalue_reference)
 #  define _CCCL_BUILTIN_IS_RVALUE_REFERENCE(...) __is_rvalue_reference(__VA_ARGS__)
-#endif // __check_builtin(is_rvalue_reference)
+#endif // __has_builtin(__is_rvalue_reference)
 
 #if __check_builtin(is_same) && !defined(_CCCL_CUDA_COMPILER_NVCC)
 #  define _CCCL_BUILTIN_IS_SAME(...) __is_same(__VA_ARGS__)
 #endif // __check_builtin(is_same)
 
-// libstdc++ defines this as a function, breaking functionality
-#if 0 // __check_builtin(is_scalar)
+// Disabled due to libstdc++ conflict
+#if 0 // __has_builtin(__is_scalar)
 #  define _CCCL_BUILTIN_IS_SCALAR(...) __is_scalar(__VA_ARGS__)
-#endif // __check_builtin(is_scalar)
+#endif // __has_builtin(__is_scalar)
 
-// libstdc++ defines this as a function, breaking functionality
-#if 0 // __check_builtin(is_signed)
+// Disabled due to libstdc++ conflict
+#if 0 // __has_builtin(__is_signed)
 #  define _CCCL_BUILTIN_IS_SIGNED(...) __is_signed(__VA_ARGS__)
-#endif // __check_builtin(is_signed)
+#endif // __has_builtin(__is_signed)
 
 #if __check_builtin(is_standard_layout) || (defined(_CCCL_COMPILER_GCC) && _CCCL_GCC_VERSION >= 40700) \
   || defined(_CCCL_COMPILER_MSVC) || defined(_CCCL_COMPILER_NVRTC)
 #  define _CCCL_BUILTIN_IS_STANDARD_LAYOUT(...) __is_standard_layout(__VA_ARGS__)
-#endif // __check_builtin(is_standard_layout)
+#endif // __check_builtin(is_standard_layout) && gcc >= 4.7
 
 #if __check_builtin(is_trivial) || (defined(_CCCL_COMPILER_GCC) && _CCCL_GCC_VERSION >= 40500) \
   || defined(_CCCL_COMPILER_MSVC) || defined(_CCCL_COMPILER_NVRTC)
 #  define _CCCL_BUILTIN_IS_TRIVIAL(...) __is_trivial(__VA_ARGS__)
-#endif // __check_builtin(is_trivial)
+#endif // __check_builtin(is_trivial) && gcc >= 4.5
 
 #if __check_builtin(is_trivially_assignable) || (defined(_CCCL_COMPILER_GCC) && _CCCL_GCC_VERSION >= 50100) \
   || defined(_CCCL_COMPILER_MSVC) || defined(_CCCL_COMPILER_NVRTC)
 #  define _CCCL_BUILTIN_IS_TRIVIALLY_ASSIGNABLE(...) __is_trivially_assignable(__VA_ARGS__)
-#endif // __check_builtin(is_trivially_assignable)
+#endif // __check_builtin(is_trivially_assignable) && gcc >= 5.1
 
 #if __check_builtin(is_trivially_constructible) || (defined(_CCCL_COMPILER_GCC) && _CCCL_GCC_VERSION >= 50100) \
   || defined(_CCCL_COMPILER_MSVC) || defined(_CCCL_COMPILER_NVRTC)
 #  define _CCCL_BUILTIN_IS_TRIVIALLY_CONSTRUCTIBLE(...) __is_trivially_constructible(__VA_ARGS__)
-#endif // __check_builtin(is_trivially_constructible)
+#endif // __check_builtin(is_trivially_constructible) && gcc >= 5.1
 
 #if __check_builtin(is_trivially_copyable) || (defined(_CCCL_COMPILER_GCC) && _CCCL_GCC_VERSION >= 50100) \
   || defined(_CCCL_COMPILER_MSVC) || defined(_CCCL_COMPILER_NVRTC)
 #  define _CCCL_BUILTIN_IS_TRIVIALLY_COPYABLE(...) __is_trivially_copyable(__VA_ARGS__)
-#endif // __check_builtin(is_trivially_copyable)
+#endif // __check_builtin(is_trivially_copyable) && gcc >= 5.1
 
 #if __check_builtin(is_trivially_destructible) || defined(_CCCL_COMPILER_MSVC)
 #  define _CCCL_BUILTIN_IS_TRIVIALLY_DESTRUCTIBLE(...) __is_trivially_destructible(__VA_ARGS__)
@@ -291,76 +340,90 @@
 #if __check_builtin(is_union) || (defined(_CCCL_COMPILER_GCC) && _CCCL_GCC_VERSION >= 40300) \
   || defined(_CCCL_COMPILER_MSVC) || defined(_CCCL_COMPILER_NVRTC)
 #  define _CCCL_BUILTIN_IS_UNION(...) __is_union(__VA_ARGS__)
-#endif // __check_builtin(is_union)
+#endif // __check_builtin(is_union) && gcc >= 4.3
 
-#if __check_builtin(is_unsigned)
+#if __check_builtin(is_unsigned) && !defined(_CCCL_CUDACC_BELOW_11_3)
 #  define _CCCL_BUILTIN_IS_UNSIGNED(...) __is_unsigned(__VA_ARGS__)
-#endif // __check_builtin(is_unsigned)
-
-#ifndef _LIBCUDACXX_USE_IS_UNSIGNED_FALLBACK
-#  if defined(_CCCL_CUDACC_BELOW_11_3)
-#    define _LIBCUDACXX_USE_IS_UNSIGNED_FALLBACK
-#  endif // _CCCL_CUDACC_BELOW_11_3
-#endif // !_LIBCUDACXX_USE_IS_UNSIGNED_FALLBACK
+#endif // __check_builtin(is_unsigned) && nvcc >= 11.3
 
-// libstdc++ defines this as a function, breaking functionality
-#if 0 // __check_builtin(is_void)
+// Disabled due to libstdc++ conflict
+#if 0 // __has_builtin(__is_void)
 #  define _CCCL_BUILTIN_IS_VOID(...) __is_void(__VA_ARGS__)
-#endif // __check_builtin(is_void)
+#endif // __has_builtin(__is_void)
+
+// Disabled due to libstdc++ conflict
+#if 0 // __has_builtin(__is_volatile)
+#  define _CCCL_BUILTIN_IS_VOLATILE(...) __is_volatile(__VA_ARGS__)
+#endif // __has_builtin(__is_volatile)
+
+#if __check_builtin(isfinite)
+#  define _CCCL_BUILTIN_ISFINITE(...) __builtin_isfinite(__VA_ARGS__)
+#endif // __check_builtin(isfinite)
+
+#if __check_builtin(isinf)
+#  define _CCCL_BUILTIN_ISINF(...) __builtin_isinf(__VA_ARGS__)
+#endif // __check_builtin(isinf)
+
+#if __check_builtin(isnan)
+#  define _CCCL_BUILTIN_ISNAN(...) __builtin_isnan(__VA_ARGS__)
+#endif // __check_builtin(isnan)
+
+#if __check_builtin(make_integer_seq) || (defined(_CCCL_COMPILER_MSVC) && _CCCL_MSVC_VERSION >= 1923)
+#  define _CCCL_BUILTIN_MAKE_INTEGER_SEQ(...) __make_integer_seq<__VA_ARGS__>
+#endif // __check_builtin(make_integer_seq)
 
 // Disabled due to libstdc++ conflict
-#if 0 // __check_builtin(make_signed)
+#if 0 // __has_builtin(__make_signed)
 #  define _CCCL_BUILTIN_MAKE_SIGNED(...) __make_signed(__VA_ARGS__)
-#endif // __check_builtin(make_signed)
+#endif // __has_builtin(__make_signed)
 
 // Disabled due to libstdc++ conflict
-#if 0 // __check_builtin(make_unsigned)
+#if 0 // __has_builtin(__make_unsigned)
 #  define _CCCL_BUILTIN_MAKE_UNSIGNED(...) __make_unsigned(__VA_ARGS__)
-#endif // __check_builtin(make_unsigned)
+#endif // __has_builtin(__make_unsigned)
 
-// Disabled due to libstdc++ conflict
-#if 0 // __check_builtin(remove_all_extents)
+#if __has_builtin(__remove_all_extents) && defined(_CCCL_CUDA_COMPILER_CLANG)
 #  define _CCCL_BUILTIN_REMOVE_ALL_EXTENTS(...) __remove_all_extents(__VA_ARGS__)
-#endif // __check_builtin(remove_all_extents)
+#endif // __has_builtin(__remove_all_extents)
 
-// Disabled due to libstdc++ conflict
-#if 0 // __check_builtin(remove_const)
+#if __has_builtin(__remove_const) && defined(_CCCL_CUDA_COMPILER_CLANG)
 #  define _CCCL_BUILTIN_REMOVE_CONST(...) __remove_const(__VA_ARGS__)
-#endif // __check_builtin(remove_const)
+#endif // __has_builtin(__remove_const)
 
-// Disabled due to libstdc++ conflict
-#if 0 // __check_builtin(remove_cv)
+#if __has_builtin(__remove_cv) && defined(_CCCL_CUDA_COMPILER_CLANG)
 #  define _CCCL_BUILTIN_REMOVE_CV(...) __remove_cv(__VA_ARGS__)
-#endif // __check_builtin(remove_cv)
+#endif // __has_builtin(__remove_cv)
 
-// Disabled due to libstdc++ conflict
-#if 0 // __check_builtin(remove_cvref)
+#if __has_builtin(__remove_cvref) && defined(_CCCL_CUDA_COMPILER_CLANG)
 #  define _CCCL_BUILTIN_REMOVE_CVREF(...) __remove_cvref(__VA_ARGS__)
-#endif // __check_builtin(remove_cvref)
+#endif // __has_builtin(__remove_cvref)
 
-// Disabled due to libstdc++ conflict
-#if 0 // __check_builtin(remove_extent)
+#if __has_builtin(__remove_extent) && defined(_CCCL_CUDA_COMPILER_CLANG)
 #  define _CCCL_BUILTIN_REMOVE_EXTENT(...) __remove_extent(__VA_ARGS__)
-#endif // __check_builtin(remove_extent)
+#endif // __has_builtin(__remove_extent)
 
-// Disabled due to libstdc++ conflict
-#if 0 // __check_builtin(remove_pointer)
+#if __has_builtin(__remove_pointer) && defined(_CCCL_CUDA_COMPILER_CLANG)
 #  define _CCCL_BUILTIN_REMOVE_POINTER(...) __remove_pointer(__VA_ARGS__)
-#endif // __check_builtin(remove_pointer)
+#endif // __has_builtin(__remove_pointer)
 
-// Disabled due to libstdc++ conflict
-#if 0 // __check_builtin(remove_reference_t)
+#if __has_builtin(__remove_reference)
+#  define _CCCL_BUILTIN_REMOVE_REFERENCE_T(...) __remove_reference(__VA_ARGS__)
+#elif __has_builtin(__remove_reference_t) && defined(_CCCL_CUDA_COMPILER_CLANG)
 #  define _CCCL_BUILTIN_REMOVE_REFERENCE_T(...) __remove_reference_t(__VA_ARGS__)
-#endif // __check_builtin(remove_reference_t)
+#endif // __has_builtin(__remove_reference_t)
 
-// Disabled due to libstdc++ conflict
-#if 0 // __check_builtin(remove_volatile)
+#if __has_builtin(__remove_volatile) && defined(_CCCL_CUDA_COMPILER_CLANG)
 #  define _CCCL_BUILTIN_REMOVE_VOLATILE(...) __remove_volatile(__VA_ARGS__)
-#endif // __check_builtin(remove_volatile)
+#endif // __has_builtin(__remove_volatile)
+
+// Versions of nvcc prior to 12.0 have trouble with pack expansion into __type_pack_element in an alias template
+#if __has_builtin(__type_pack_element) && !defined(_CCCL_CUDACC_BELOW_12_0)
+#  define _CCCL_BUILTIN_TYPE_PACK_ELEMENT(...) __type_pack_element<__VA_ARGS__>
+#endif // __has_builtin(__type_pack_element)
 
 #if __check_builtin(underlying_type) || (defined(_CCCL_COMPILER_GCC) && _CCCL_GCC_VERSION >= 40700) \
   || defined(_CCCL_COMPILER_MSVC) || defined(_CCCL_COMPILER_NVRTC)
 #  define _CCCL_BUILTIN_UNDERLYING_TYPE(...) __underlying_type(__VA_ARGS__)
-#endif // __check_builtin(underlying_type)
+#endif // __check_builtin(underlying_type) && gcc >= 4.7
 
 #endif // __CCCL_BUILTIN_H
diff --git a/libcudacxx/include/cuda/std/__memory/addressof.h b/libcudacxx/include/cuda/std/__memory/addressof.h
index d6e7cee116..2a04ff0e4d 100644
--- a/libcudacxx/include/cuda/std/__memory/addressof.h
+++ b/libcudacxx/include/cuda/std/__memory/addressof.h
@@ -30,7 +30,7 @@ _LIBCUDACXX_BEGIN_NAMESPACE_STD
 template <class _Tp>
 _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 _LIBCUDACXX_NO_CFI _Tp* addressof(_Tp& __x) noexcept
 {
-  return __builtin_addressof(__x);
+  return _CCCL_BUILTIN_ADDRESSOF(__x);
 }
 
 #else
diff --git a/libcudacxx/include/cuda/std/__new/allocate.h b/libcudacxx/include/cuda/std/__new/allocate.h
index 9081dfcf81..b37e934872 100644
--- a/libcudacxx/include/cuda/std/__new/allocate.h
+++ b/libcudacxx/include/cuda/std/__new/allocate.h
@@ -60,22 +60,22 @@ template <class... _Args>
 _LIBCUDACXX_HIDE_FROM_ABI void* __libcpp_operator_new(_Args... __args)
 {
   // Those builtins are not usable on device and the tests crash when using them
-#if 0 && __has_builtin(__builtin_operator_new) && __has_builtin(__builtin_operator_delete)
-  return __builtin_operator_new(__args...);
-#else // ^^^ use builtin ^^^ / vvv no builtin
+#if defined(_CCCL_BUILTIN_OPERATOR_NEW)
+  return _CCCL_BUILTIN_OPERATOR_NEW(__args...);
+#else // ^^^ _CCCL_BUILTIN_OPERATOR_NEW ^^^ / vvv !_CCCL_BUILTIN_OPERATOR_NEW vvv
   return ::operator new(__args...);
-#endif // !__builtin_operator_new || !__builtin_operator_delete
+#endif // !_CCCL_BUILTIN_OPERATOR_NEW
 }
 
 template <class... _Args>
 _LIBCUDACXX_HIDE_FROM_ABI void __libcpp_operator_delete(_Args... __args)
 {
   // Those builtins are not usable on device and the tests crash when using them
-#if 0 && __has_builtin(__builtin_operator_new) && __has_builtin(__builtin_operator_delete)
-  __builtin_operator_delete(__args...);
-#else // ^^^ use builtin ^^^ / vvv no builtin
+#if defined(_CCCL_BUILTIN_OPERATOR_DELETE)
+  _CCCL_BUILTIN_OPERATOR_DELETE(__args...);
+#else // ^^^ _CCCL_BUILTIN_OPERATOR_DELETE ^^^ / vvv !_CCCL_BUILTIN_OPERATOR_DELETE vvv
   ::operator delete(__args...);
-#endif // !__builtin_operator_new || !__builtin_operator_delete
+#endif // !_CCCL_BUILTIN_OPERATOR_DELETE
 }
 
 _LIBCUDACXX_HIDE_FROM_ABI void* __libcpp_allocate(size_t __size, size_t __align)
diff --git a/libcudacxx/include/cuda/std/__new/launder.h b/libcudacxx/include/cuda/std/__new/launder.h
index 4d6e81e427..dbce092543 100644
--- a/libcudacxx/include/cuda/std/__new/launder.h
+++ b/libcudacxx/include/cuda/std/__new/launder.h
@@ -37,7 +37,7 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 _Tp* launder(_Tp
   return _CCCL_BUILTIN_LAUNDER(__p);
 #else
   return __p;
-#endif // defined(_CCCL_BUILTIN_LAUNDER)
+#endif // _CCCL_BUILTIN_LAUNDER
 }
 
 _LIBCUDACXX_END_NAMESPACE_STD
diff --git a/libcudacxx/include/cuda/std/__type_traits/add_lvalue_reference.h b/libcudacxx/include/cuda/std/__type_traits/add_lvalue_reference.h
index c7fe3b3c30..7e729d5c5b 100644
--- a/libcudacxx/include/cuda/std/__type_traits/add_lvalue_reference.h
+++ b/libcudacxx/include/cuda/std/__type_traits/add_lvalue_reference.h
@@ -24,12 +24,12 @@
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
-#if defined(_LIBCUDACXX_ADD_LVALUE_REFERENCE) && !defined(_LIBCUDACXX_USE_ADD_LVALUE_REFERENCE_FALLBACK)
+#if defined(_CCCL_BUILTIN_ADD_LVALUE_REFERENCE) && !defined(_LIBCUDACXX_USE_ADD_LVALUE_REFERENCE_FALLBACK)
 
 template <class _Tp>
-using __add_lvalue_reference_t = _LIBCUDACXX_ADD_LVALUE_REFERENCE(_Tp);
+using __add_lvalue_reference_t = _CCCL_BUILTIN_ADD_LVALUE_REFERENCE(_Tp);
 
-#else
+#else // ^^^ _CCCL_BUILTIN_ADD_LVALUE_REFERENCE ^^^ / vvv !_CCCL_BUILTIN_ADD_LVALUE_REFERENCE vvv
 
 template <class _Tp, bool = __libcpp_is_referenceable<_Tp>::value>
 struct __add_lvalue_reference_impl
@@ -45,7 +45,7 @@ struct __add_lvalue_reference_impl<_Tp, true>
 template <class _Tp>
 using __add_lvalue_reference_t = typename __add_lvalue_reference_impl<_Tp>::type;
 
-#endif // defined(_LIBCUDACXX_ADD_LVALUE_REFERENCE) && !defined(_LIBCUDACXX_USE_ADD_LVALUE_REFERENCE_FALLBACK)
+#endif // !_CCCL_BUILTIN_ADD_LVALUE_REFERENCE
 
 template <class _Tp>
 struct add_lvalue_reference
diff --git a/libcudacxx/include/cuda/std/__type_traits/add_pointer.h b/libcudacxx/include/cuda/std/__type_traits/add_pointer.h
index 856c48beee..a830ad3b30 100644
--- a/libcudacxx/include/cuda/std/__type_traits/add_pointer.h
+++ b/libcudacxx/include/cuda/std/__type_traits/add_pointer.h
@@ -28,12 +28,12 @@
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
-#if defined(_LIBCUDACXX_ADD_POINTER) && !defined(_LIBCUDACXX_USE_ADD_POINTER_FALLBACK)
+#if defined(_CCCL_BUILTIN_ADD_POINTER) && !defined(_LIBCUDACXX_USE_ADD_POINTER_FALLBACK)
 
 template <class _Tp>
-using __add_pointer_t = _LIBCUDACXX_ADD_POINTER(_Tp);
+using __add_pointer_t = _CCCL_BUILTIN_ADD_POINTER(_Tp);
 
-#else
+#else // ^^^ _CCCL_BUILTIN_ADD_POINTER ^^^ / vvv !_CCCL_BUILTIN_ADD_POINTER vvv
 template <class _Tp, bool = __libcpp_is_referenceable<_Tp>::value || is_void<_Tp>::value>
 struct __add_pointer_impl
 {
@@ -48,7 +48,7 @@ struct __add_pointer_impl<_Tp, false>
 template <class _Tp>
 using __add_pointer_t = typename __add_pointer_impl<_Tp>::type;
 
-#endif // defined(_LIBCUDACXX_ADD_POINTER) && !defined(_LIBCUDACXX_USE_ADD_POINTER_FALLBACK)
+#endif // !_CCCL_BUILTIN_ADD_POINTER
 
 template <class _Tp>
 struct add_pointer
diff --git a/libcudacxx/include/cuda/std/__type_traits/add_rvalue_reference.h b/libcudacxx/include/cuda/std/__type_traits/add_rvalue_reference.h
index 8e3f79e856..9796ce1f22 100644
--- a/libcudacxx/include/cuda/std/__type_traits/add_rvalue_reference.h
+++ b/libcudacxx/include/cuda/std/__type_traits/add_rvalue_reference.h
@@ -24,12 +24,12 @@
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
-#if defined(_LIBCUDACXX_ADD_RVALUE_REFERENCE) && !defined(_LIBCUDACXX_USE_ADD_RVALUE_REFERENCE_FALLBACK)
+#if defined(_CCCL_BUILTIN_ADD_RVALUE_REFERENCE) && !defined(_LIBCUDACXX_USE_ADD_RVALUE_REFERENCE_FALLBACK)
 
 template <class _Tp>
-using __add_rvalue_reference_t = _LIBCUDACXX_ADD_RVALUE_REFERENCE(_Tp);
+using __add_rvalue_reference_t = _CCCL_BUILTIN_ADD_RVALUE_REFERENCE(_Tp);
 
-#else
+#else // ^^^ _CCCL_BUILTIN_ADD_RVALUE_REFERENCE ^^^ / vvv !_CCCL_BUILTIN_ADD_RVALUE_REFERENCE vvv
 
 template <class _Tp, bool = __libcpp_is_referenceable<_Tp>::value>
 struct __add_rvalue_reference_impl
@@ -45,7 +45,7 @@ struct __add_rvalue_reference_impl<_Tp, true>
 template <class _Tp>
 using __add_rvalue_reference_t = typename __add_rvalue_reference_impl<_Tp>::type;
 
-#endif // defined(_LIBCUDACXX_ADD_RVALUE_REFERENCE) && !defined(_LIBCUDACXX_USE_ADD_RVALUE_REFERENCE_FALLBACK)
+#endif // _CCCL_BUILTIN_ADD_RVALUE_REFERENCE
 
 template <class _Tp>
 struct add_rvalue_reference
diff --git a/libcudacxx/include/cuda/std/__type_traits/extent.h b/libcudacxx/include/cuda/std/__type_traits/extent.h
index 75bf537297..8f81bbaf32 100644
--- a/libcudacxx/include/cuda/std/__type_traits/extent.h
+++ b/libcudacxx/include/cuda/std/__type_traits/extent.h
@@ -26,18 +26,18 @@
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
-#if defined(_LIBCUDACXX_ARRAY_EXTENT) && !defined(_LIBCUDACXX_USE_ARRAY_EXTENT_FALLBACK)
+#if defined(_CCCL_BUILTIN_ARRAY_EXTENT) && !defined(_LIBCUDACXX_USE_ARRAY_EXTENT_FALLBACK)
 
 template <class _Tp, size_t _Dim = 0>
-struct _CCCL_TYPE_VISIBILITY_DEFAULT extent : integral_constant<size_t, _LIBCUDACXX_ARRAY_EXTENT(_Tp, _Dim)>
+struct _CCCL_TYPE_VISIBILITY_DEFAULT extent : integral_constant<size_t, _CCCL_BUILTIN_ARRAY_EXTENT(_Tp, _Dim)>
 {};
 
 #  if _CCCL_STD_VER > 2011 && !defined(_LIBCUDACXX_HAS_NO_VARIABLE_TEMPLATES)
 template <class _Tp, unsigned _Ip = 0>
-_LIBCUDACXX_INLINE_VAR constexpr size_t extent_v = _LIBCUDACXX_ARRAY_EXTENT(_Tp, _Ip);
+_LIBCUDACXX_INLINE_VAR constexpr size_t extent_v = _CCCL_BUILTIN_ARRAY_EXTENT(_Tp, _Ip);
 #  endif
 
-#else
+#else // ^^^ _CCCL_BUILTIN_ARRAY_EXTENT ^^^ / vvv !_CCCL_BUILTIN_ARRAY_EXTENT vvv
 
 template <class _Tp, unsigned _Ip = 0>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT extent : public integral_constant<size_t, 0>
@@ -61,7 +61,7 @@ template <class _Tp, unsigned _Ip = 0>
 _LIBCUDACXX_INLINE_VAR constexpr size_t extent_v = extent<_Tp, _Ip>::value;
 #  endif
 
-#endif // defined(_LIBCUDACXX_ARRAY_EXTENT) && !defined(_LIBCUDACXX_USE_ARRAY_EXTENT_FALLBACK)
+#endif // !_CCCL_BUILTIN_ARRAY_EXTENT
 
 _LIBCUDACXX_END_NAMESPACE_STD
 
diff --git a/libcudacxx/include/cuda/std/__type_traits/is_compound.h b/libcudacxx/include/cuda/std/__type_traits/is_compound.h
index f400a99567..356c2d25b9 100644
--- a/libcudacxx/include/cuda/std/__type_traits/is_compound.h
+++ b/libcudacxx/include/cuda/std/__type_traits/is_compound.h
@@ -25,18 +25,18 @@
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
-#if defined(_LIBCUDACXX_IS_COMPOUND) && !defined(_LIBCUDACXX_USE_IS_COMPOUND_FALLBACK)
+#if defined(_CCCL_BUILTIN_IS_COMPOUND) && !defined(_LIBCUDACXX_USE_IS_COMPOUND_FALLBACK)
 
 template <class _Tp>
-struct _CCCL_TYPE_VISIBILITY_DEFAULT is_compound : public integral_constant<bool, _LIBCUDACXX_IS_COMPOUND(_Tp)>
+struct _CCCL_TYPE_VISIBILITY_DEFAULT is_compound : public integral_constant<bool, _CCCL_BUILTIN_IS_COMPOUND(_Tp)>
 {};
 
 #  if _CCCL_STD_VER > 2011 && !defined(_LIBCUDACXX_HAS_NO_VARIABLE_TEMPLATES)
 template <class _Tp>
-_LIBCUDACXX_INLINE_VAR constexpr bool is_compound_v = _LIBCUDACXX_IS_COMPOUND(_Tp);
+_LIBCUDACXX_INLINE_VAR constexpr bool is_compound_v = _CCCL_BUILTIN_IS_COMPOUND(_Tp);
 #  endif
 
-#else
+#else // ^^^ _CCCL_BUILTIN_IS_COMPOUND ^^^ / vvv !_CCCL_BUILTIN_IS_COMPOUND vvv
 
 template <class _Tp>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT is_compound : public integral_constant<bool, !is_fundamental<_Tp>::value>
@@ -47,7 +47,7 @@ template <class _Tp>
 _LIBCUDACXX_INLINE_VAR constexpr bool is_compound_v = is_compound<_Tp>::value;
 #  endif
 
-#endif // defined(_LIBCUDACXX_IS_COMPOUND) && !defined(_LIBCUDACXX_USE_IS_COMPOUND_FALLBACK)
+#endif // !_CCCL_BUILTIN_IS_COMPOUND
 
 _LIBCUDACXX_END_NAMESPACE_STD
 
diff --git a/libcudacxx/include/cuda/std/__type_traits/is_const.h b/libcudacxx/include/cuda/std/__type_traits/is_const.h
index 7a4d6702bc..94d669e390 100644
--- a/libcudacxx/include/cuda/std/__type_traits/is_const.h
+++ b/libcudacxx/include/cuda/std/__type_traits/is_const.h
@@ -24,18 +24,18 @@
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
-#if defined(_LIBCUDACXX_IS_CONST) && !defined(_LIBCUDACXX_USE_IS_CONST_FALLBACK)
+#if defined(_CCCL_BUILTIN_IS_CONST) && !defined(_LIBCUDACXX_USE_IS_CONST_FALLBACK)
 
 template <class _Tp>
-struct _CCCL_TYPE_VISIBILITY_DEFAULT is_const : public integral_constant<bool, _LIBCUDACXX_IS_CONST(_Tp)>
+struct _CCCL_TYPE_VISIBILITY_DEFAULT is_const : public integral_constant<bool, _CCCL_BUILTIN_IS_CONST(_Tp)>
 {};
 
 #  if _CCCL_STD_VER > 2011 && !defined(_LIBCUDACXX_HAS_NO_VARIABLE_TEMPLATES)
 template <class _Tp>
-_LIBCUDACXX_INLINE_VAR constexpr bool is_const_v = _LIBCUDACXX_IS_CONST(_Tp);
+_LIBCUDACXX_INLINE_VAR constexpr bool is_const_v = _CCCL_BUILTIN_IS_CONST(_Tp);
 #  endif
 
-#else
+#else // ^^^ _CCCL_BUILTIN_IS_CONST ^^^ / vvv !_CCCL_BUILTIN_IS_CONST vvv
 
 template <class _Tp>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT is_const : public false_type
@@ -49,7 +49,7 @@ template <class _Tp>
 _LIBCUDACXX_INLINE_VAR constexpr bool is_const_v = is_const<_Tp>::value;
 #  endif
 
-#endif // defined(_LIBCUDACXX_IS_CONST) && !defined(_LIBCUDACXX_USE_IS_CONST_FALLBACK)
+#endif // !_CCCL_BUILTIN_IS_CONST
 
 _LIBCUDACXX_END_NAMESPACE_STD
 
diff --git a/libcudacxx/include/cuda/std/__type_traits/is_fundamental.h b/libcudacxx/include/cuda/std/__type_traits/is_fundamental.h
index dd75d9c50c..0e3b2060f7 100644
--- a/libcudacxx/include/cuda/std/__type_traits/is_fundamental.h
+++ b/libcudacxx/include/cuda/std/__type_traits/is_fundamental.h
@@ -27,18 +27,18 @@
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
-#if defined(_LIBCUDACXX_IS_FUNDAMENTAL) && !defined(_LIBCUDACXX_USE_IS_FUNDAMENTAL_FALLBACK)
+#if defined(_CCCL_BUILTIN_IS_FUNDAMENTAL) && !defined(_LIBCUDACXX_USE_IS_FUNDAMENTAL_FALLBACK)
 
 template <class _Tp>
-struct _CCCL_TYPE_VISIBILITY_DEFAULT is_fundamental : public integral_constant<bool, _LIBCUDACXX_IS_FUNDAMENTAL(_Tp)>
+struct _CCCL_TYPE_VISIBILITY_DEFAULT is_fundamental : public integral_constant<bool, _CCCL_BUILTIN_IS_FUNDAMENTAL(_Tp)>
 {};
 
 #  if _CCCL_STD_VER > 2011 && !defined(_LIBCUDACXX_HAS_NO_VARIABLE_TEMPLATES)
 template <class _Tp>
-_LIBCUDACXX_INLINE_VAR constexpr bool is_fundamental_v = _LIBCUDACXX_IS_FUNDAMENTAL(_Tp);
+_LIBCUDACXX_INLINE_VAR constexpr bool is_fundamental_v = _CCCL_BUILTIN_IS_FUNDAMENTAL(_Tp);
 #  endif
 
-#else
+#else // ^^^ _CCCL_BUILTIN_IS_FUNDAMENTAL ^^^ / vvv !_CCCL_BUILTIN_IS_FUNDAMENTAL vvv
 
 template <class _Tp>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT is_fundamental
@@ -50,7 +50,7 @@ template <class _Tp>
 _LIBCUDACXX_INLINE_VAR constexpr bool is_fundamental_v = is_fundamental<_Tp>::value;
 #  endif
 
-#endif // defined(_LIBCUDACXX_IS_FUNDAMENTAL) && !defined(_LIBCUDACXX_USE_IS_FUNDAMENTAL_FALLBACK)
+#endif // !_CCCL_BUILTIN_IS_FUNDAMENTAL
 
 _LIBCUDACXX_END_NAMESPACE_STD
 
diff --git a/libcudacxx/include/cuda/std/__type_traits/is_integral.h b/libcudacxx/include/cuda/std/__type_traits/is_integral.h
index 5100215119..89562a0be8 100644
--- a/libcudacxx/include/cuda/std/__type_traits/is_integral.h
+++ b/libcudacxx/include/cuda/std/__type_traits/is_integral.h
@@ -25,18 +25,18 @@
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
-#if defined(_LIBCUDACXX_IS_INTEGRAL) && !defined(_LIBCUDACXX_USE_IS_INTEGRAL_FALLBACK)
+#if defined(_CCCL_BUILTIN_IS_INTEGRAL) && !defined(_LIBCUDACXX_USE_IS_INTEGRAL_FALLBACK)
 
 template <class _Tp>
-struct _CCCL_TYPE_VISIBILITY_DEFAULT is_integral : public integral_constant<bool, _LIBCUDACXX_IS_INTEGRAL(_Tp)>
+struct _CCCL_TYPE_VISIBILITY_DEFAULT is_integral : public integral_constant<bool, _CCCL_BUILTIN_IS_INTEGRAL(_Tp)>
 {};
 
 #  if _CCCL_STD_VER > 2011 && !defined(_LIBCUDACXX_HAS_NO_VARIABLE_TEMPLATES)
 template <class _Tp>
-_LIBCUDACXX_INLINE_VAR constexpr bool is_integral_v = _LIBCUDACXX_IS_INTEGRAL(_Tp);
+_LIBCUDACXX_INLINE_VAR constexpr bool is_integral_v = _CCCL_BUILTIN_IS_INTEGRAL(_Tp);
 #  endif
 
-#else
+#else // ^^^ _CCCL_BUILTIN_IS_INTEGRAL ^^^ / vvv !_CCCL_BUILTIN_IS_INTEGRAL vvv
 
 template <class _Tp>
 struct __libcpp_is_integral : public false_type
@@ -112,7 +112,7 @@ template <class _Tp>
 _LIBCUDACXX_INLINE_VAR constexpr bool is_integral_v = is_integral<_Tp>::value;
 #  endif
 
-#endif // defined(_LIBCUDACXX_IS_INTEGRAL) && !defined(_LIBCUDACXX_USE_IS_INTEGRAL_FALLBACK)
+#endif // !_CCCL_BUILTIN_IS_INTEGRAL
 
 _LIBCUDACXX_END_NAMESPACE_STD
 
diff --git a/libcudacxx/include/cuda/std/__type_traits/is_member_function_pointer.h b/libcudacxx/include/cuda/std/__type_traits/is_member_function_pointer.h
index f6963b1d96..4b2b26b882 100644
--- a/libcudacxx/include/cuda/std/__type_traits/is_member_function_pointer.h
+++ b/libcudacxx/include/cuda/std/__type_traits/is_member_function_pointer.h
@@ -27,20 +27,6 @@
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
-#if defined(_LIBCUDACXX_IS_MEMBER_FUNCTION_POINTER) && !defined(_LIBCUDACXX_USE_IS_MEMBER_FUNCTION_POINTER_FALLBACK)
-
-template <class _Tp>
-struct _CCCL_TYPE_VISIBILITY_DEFAULT is_member_function_pointer
-    : public integral_constant<bool, _LIBCUDACXX_IS_MEMBER_FUNCTION_POINTER(_Tp)>
-{};
-
-#  if _CCCL_STD_VER > 2011 && !defined(_LIBCUDACXX_HAS_NO_VARIABLE_TEMPLATES)
-template <class _Tp>
-_LIBCUDACXX_INLINE_VAR constexpr bool is_member_function_pointer_v = _LIBCUDACXX_IS_MEMBER_FUNCTION_POINTER(_Tp);
-#  endif
-
-#else
-
 template <class _Tp>
 struct __libcpp_is_member_pointer
 {
@@ -62,6 +48,20 @@ struct __libcpp_is_member_pointer<_Tp _Up::*>
   };
 };
 
+#if defined(_CCCL_BUILTIN_IS_MEMBER_FUNCTION_POINTER) && !defined(_LIBCUDACXX_USE_IS_MEMBER_FUNCTION_POINTER_FALLBACK)
+
+template <class _Tp>
+struct _CCCL_TYPE_VISIBILITY_DEFAULT is_member_function_pointer
+    : public integral_constant<bool, _CCCL_BUILTIN_IS_MEMBER_FUNCTION_POINTER(_Tp)>
+{};
+
+#  if _CCCL_STD_VER > 2011 && !defined(_LIBCUDACXX_HAS_NO_VARIABLE_TEMPLATES)
+template <class _Tp>
+_LIBCUDACXX_INLINE_VAR constexpr bool is_member_function_pointer_v = _CCCL_BUILTIN_IS_MEMBER_FUNCTION_POINTER(_Tp);
+#  endif
+
+#else // ^^^ _CCCL_BUILTIN_IS_MEMBER_FUNCTION_POINTER ^^^ / vvv !_CCCL_BUILTIN_IS_MEMBER_FUNCTION_POINTER vvv
+
 template <class _Tp>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT is_member_function_pointer
     : public integral_constant<bool, __libcpp_is_member_pointer<__remove_cv_t<_Tp>>::__is_func>
@@ -72,8 +72,7 @@ template <class _Tp>
 _LIBCUDACXX_INLINE_VAR constexpr bool is_member_function_pointer_v = is_member_function_pointer<_Tp>::value;
 #  endif
 
-#endif // defined(_LIBCUDACXX_IS_MEMBER_FUNCTION_POINTER) &&
-       // !defined(_LIBCUDACXX_USE_IS_MEMBER_FUNCTION_POINTER_FALLBACK)
+#endif // !_CCCL_BUILTIN_IS_MEMBER_FUNCTION_POINTER
 
 _LIBCUDACXX_END_NAMESPACE_STD
 
diff --git a/libcudacxx/include/cuda/std/__type_traits/is_member_object_pointer.h b/libcudacxx/include/cuda/std/__type_traits/is_member_object_pointer.h
index 93860926f1..391df40c14 100644
--- a/libcudacxx/include/cuda/std/__type_traits/is_member_object_pointer.h
+++ b/libcudacxx/include/cuda/std/__type_traits/is_member_object_pointer.h
@@ -26,19 +26,19 @@
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
-#if defined(_LIBCUDACXX_IS_MEMBER_OBJECT_POINTER) && !defined(_LIBCUDACXX_USE_IS_MEMBER_OBJECT_POINTER_FALLBACK)
+#if defined(_CCCL_BUILTIN_IS_MEMBER_OBJECT_POINTER) && !defined(_LIBCUDACXX_USE_IS_MEMBER_OBJECT_POINTER_FALLBACK)
 
 template <class _Tp>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT is_member_object_pointer
-    : public integral_constant<bool, _LIBCUDACXX_IS_MEMBER_OBJECT_POINTER(_Tp)>
+    : public integral_constant<bool, _CCCL_BUILTIN_IS_MEMBER_OBJECT_POINTER(_Tp)>
 {};
 
 #  if _CCCL_STD_VER > 2011 && !defined(_LIBCUDACXX_HAS_NO_VARIABLE_TEMPLATES)
 template <class _Tp>
-_LIBCUDACXX_INLINE_VAR constexpr bool is_member_object_pointer_v = _LIBCUDACXX_IS_MEMBER_OBJECT_POINTER(_Tp);
+_LIBCUDACXX_INLINE_VAR constexpr bool is_member_object_pointer_v = _CCCL_BUILTIN_IS_MEMBER_OBJECT_POINTER(_Tp);
 #  endif
 
-#else
+#else // ^^^ _CCCL_BUILTIN_IS_MEMBER_OBJECT_POINTER ^^^ / vvv !_CCCL_BUILTIN_IS_MEMBER_OBJECT_POINTER vvv
 
 template <class _Tp>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT is_member_object_pointer
@@ -50,7 +50,7 @@ template <class _Tp>
 _LIBCUDACXX_INLINE_VAR constexpr bool is_member_object_pointer_v = is_member_object_pointer<_Tp>::value;
 #  endif
 
-#endif // defined(_LIBCUDACXX_IS_MEMBER_OBJECT_POINTER) && !defined(_LIBCUDACXX_USE_IS_MEMBER_OBJECT_POINTER_FALLBACK)
+#endif // !_CCCL_BUILTIN_IS_MEMBER_OBJECT_POINTER
 
 _LIBCUDACXX_END_NAMESPACE_STD
 
diff --git a/libcudacxx/include/cuda/std/__type_traits/is_member_pointer.h b/libcudacxx/include/cuda/std/__type_traits/is_member_pointer.h
index 9289061ddf..6d6b5448d2 100644
--- a/libcudacxx/include/cuda/std/__type_traits/is_member_pointer.h
+++ b/libcudacxx/include/cuda/std/__type_traits/is_member_pointer.h
@@ -26,19 +26,19 @@
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
-#if defined(_LIBCUDACXX_IS_MEMBER_POINTER) && !defined(_LIBCUDACXX_USE_IS_MEMBER_POINTER_FALLBACK)
+#if defined(_CCCL_BUILTIN_IS_MEMBER_POINTER) && !defined(_LIBCUDACXX_USE_IS_MEMBER_POINTER_FALLBACK)
 
 template <class _Tp>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT is_member_pointer
-    : public integral_constant<bool, _LIBCUDACXX_IS_MEMBER_POINTER(_Tp)>
+    : public integral_constant<bool, _CCCL_BUILTIN_IS_MEMBER_POINTER(_Tp)>
 {};
 
 #  if _CCCL_STD_VER > 2011 && !defined(_LIBCUDACXX_HAS_NO_VARIABLE_TEMPLATES)
 template <class _Tp>
-_LIBCUDACXX_INLINE_VAR constexpr bool is_member_pointer_v = _LIBCUDACXX_IS_MEMBER_POINTER(_Tp);
+_LIBCUDACXX_INLINE_VAR constexpr bool is_member_pointer_v = _CCCL_BUILTIN_IS_MEMBER_POINTER(_Tp);
 #  endif
 
-#else
+#else // ^^^ _CCCL_BUILTIN_IS_MEMBER_POINTER ^^^ / vvv !_CCCL_BUILTIN_IS_MEMBER_POINTER vvv
 
 template <class _Tp>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT is_member_pointer
@@ -50,7 +50,7 @@ template <class _Tp>
 _LIBCUDACXX_INLINE_VAR constexpr bool is_member_pointer_v = is_member_pointer<_Tp>::value;
 #  endif
 
-#endif // defined(_LIBCUDACXX_IS_MEMBER_POINTER) && !defined(_LIBCUDACXX_USE_IS_MEMBER_POINTER_FALLBACK)
+#endif // !_CCCL_BUILTIN_IS_MEMBER_POINTER
 
 _LIBCUDACXX_END_NAMESPACE_STD
 
diff --git a/libcudacxx/include/cuda/std/__type_traits/is_volatile.h b/libcudacxx/include/cuda/std/__type_traits/is_volatile.h
index cd02b7c404..2b9f38e63e 100644
--- a/libcudacxx/include/cuda/std/__type_traits/is_volatile.h
+++ b/libcudacxx/include/cuda/std/__type_traits/is_volatile.h
@@ -24,18 +24,18 @@
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
-#if defined(_LIBCUDACXX_IS_VOLATILE) && !defined(_LIBCUDACXX_USE_IS_VOLATILE_FALLBACK)
+#if defined(_CCCL_BUILTIN_IS_VOLATILE) && !defined(_LIBCUDACXX_USE_IS_VOLATILE_FALLBACK)
 
 template <class _Tp>
-struct _CCCL_TYPE_VISIBILITY_DEFAULT is_volatile : : public integral_constant<bool, _LIBCUDACXX_IS_VOLATILE(_Tp)>
+struct _CCCL_TYPE_VISIBILITY_DEFAULT is_volatile : : public integral_constant<bool, _CCCL_BUILTIN_IS_VOLATILE(_Tp)>
 {};
 
 #  if _CCCL_STD_VER > 2011 && !defined(_LIBCUDACXX_HAS_NO_VARIABLE_TEMPLATES)
 template <class _Tp>
-_LIBCUDACXX_INLINE_VAR constexpr bool is_volatile_v = _LIBCUDACXX_IS_VOLATILE(_Tp);
+_LIBCUDACXX_INLINE_VAR constexpr bool is_volatile_v = _CCCL_BUILTIN_IS_VOLATILE(_Tp);
 #  endif
 
-#else
+#else // ^^^ _CCCL_BUILTIN_IS_VOLATILE ^^^ / vvv !_CCCL_BUILTIN_IS_VOLATILE vvv
 
 template <class _Tp>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT is_volatile : public false_type
@@ -49,7 +49,7 @@ template <class _Tp>
 _LIBCUDACXX_INLINE_VAR constexpr bool is_volatile_v = is_volatile<_Tp>::value;
 #  endif
 
-#endif // defined(_LIBCUDACXX_IS_VOLATILE) && !defined(_LIBCUDACXX_USE_IS_VOLATILE_FALLBACK)
+#endif // !_CCCL_BUILTIN_IS_VOLATILE
 
 _LIBCUDACXX_END_NAMESPACE_STD
 
diff --git a/libcudacxx/include/cuda/std/__type_traits/remove_reference.h b/libcudacxx/include/cuda/std/__type_traits/remove_reference.h
index f1b4d40a06..efc025de64 100644
--- a/libcudacxx/include/cuda/std/__type_traits/remove_reference.h
+++ b/libcudacxx/include/cuda/std/__type_traits/remove_reference.h
@@ -31,10 +31,16 @@ struct remove_reference
   using type _LIBCUDACXX_NODEBUG_TYPE = _CCCL_BUILTIN_REMOVE_REFERENCE_T(_Tp);
 };
 
+#  if defined(_CCCL_COMPILER_GCC)
+// error: use of built-in trait in function signature; use library traits instead
+template <class _Tp>
+using __libcpp_remove_reference_t = typename remove_reference<_Tp>::type;
+#  else // ^^^ _CCCL_COMPILER_GCC ^^^^/  vvv !_CCCL_COMPILER_GCC
 template <class _Tp>
 using __libcpp_remove_reference_t = _CCCL_BUILTIN_REMOVE_REFERENCE_T(_Tp);
+#  endif // !_CCCL_COMPILER_GCC
 
-#else
+#else // ^^^ _CCCL_BUILTIN_REMOVE_REFERENCE_T ^^^ / vvv !_CCCL_BUILTIN_REMOVE_REFERENCE_T vvv
 
 template <class _Tp>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT remove_reference
@@ -55,7 +61,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT remove_reference<_Tp&&>
 template <class _Tp>
 using __libcpp_remove_reference_t = typename remove_reference<_Tp>::type;
 
-#endif // defined(_CCCL_BUILTIN_REMOVE_REFERENCE_T) && !defined(_LIBCUDACXX_USE_REMOVE_REFERENCE_T_FALLBACK)
+#endif // !_CCCL_BUILTIN_REMOVE_REFERENCE_T
 
 #if _CCCL_STD_VER > 2011
 template <class _Tp>
diff --git a/libcudacxx/include/cuda/std/__type_traits/type_list.h b/libcudacxx/include/cuda/std/__type_traits/type_list.h
index 1ebdecf4d6..0be25bf6a3 100644
--- a/libcudacxx/include/cuda/std/__type_traits/type_list.h
+++ b/libcudacxx/include/cuda/std/__type_traits/type_list.h
@@ -269,7 +269,7 @@ using __type_index = _Ts...[_Ip::value];
 // Versions of nvcc prior to 12.0 have trouble with pack expansion into
 // __type_pack_element in an alias template, so we use the fall-back
 // implementation instead.
-#  elif __has_builtin(__type_pack_element) && !defined(_CCCL_CUDACC_BELOW_12_0)
+#  elif defined(_CCCL_BUILTIN_TYPE_PACK_ELEMENT)
 
 namespace __detail
 {
@@ -279,7 +279,7 @@ template <size_t _Ip>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT __type_index_fn
 {
   template <class... _Ts>
-  using __call _LIBCUDACXX_NODEBUG_TYPE = __type_pack_element<_Ip, _Ts...>;
+  using __call _LIBCUDACXX_NODEBUG_TYPE = _CCCL_BUILTIN_TYPE_PACK_ELEMENT(_Ip, _Ts...);
 };
 } // namespace __detail
 
@@ -289,7 +289,7 @@ using __type_index_c = __type_call<__detail::__type_index_fn<_Ip>, _Ts...>;
 template <class _Ip, class... _Ts>
 using __type_index = __type_call<__detail::__type_index_fn<_Ip::value>, _Ts...>;
 
-#  else // ^^^ __has_builtin(__type_pack_element) ^^^ / vvv !__has_builtin(__type_pack_element) vvv
+#  else // ^^^ _CCCL_BUILTIN_TYPE_PACK_ELEMENT ^^^ / vvv !_CCCL_BUILTIN_TYPE_PACK_ELEMENT vvv
 
 // Fallback implementation for __type_index uses multiple inheritance. See:
 // https://ldionne.com/2015/11/29/efficient-parameter-pack-indexing/
@@ -360,7 +360,7 @@ using __type_index = __type_call<__detail::__type_index_select_fn<(_Ip::value <
 template <size_t _Ip, class... _Ts>
 using __type_index_c = __type_index<integral_constant<size_t, _Ip>, _Ts...>;
 
-#  endif // !__has_builtin(__type_pack_element)
+#  endif // !_CCCL_BUILTIN_TYPE_PACK_ELEMENT
 
 namespace __detail
 {
diff --git a/libcudacxx/include/cuda/std/__utility/integer_sequence.h b/libcudacxx/include/cuda/std/__utility/integer_sequence.h
index a4feb96432..4004a38c64 100644
--- a/libcudacxx/include/cuda/std/__utility/integer_sequence.h
+++ b/libcudacxx/include/cuda/std/__utility/integer_sequence.h
@@ -39,20 +39,19 @@ struct __integer_sequence
   using __to_tuple_indices _LIBCUDACXX_NODEBUG_TYPE = __tuple_indices<(_Values + _Sp)...>;
 };
 
-#if defined(_LIBCUDACXX_HAS_MAKE_INTEGER_SEQ)
+#if defined(_CCCL_BUILTIN_MAKE_INTEGER_SEQ)
 
 template <size_t _Ep, size_t _Sp>
 using __make_indices_imp _LIBCUDACXX_NODEBUG_TYPE =
-  typename __make_integer_seq<__integer_sequence, size_t, _Ep - _Sp>::template __to_tuple_indices<_Sp>;
+  typename _CCCL_BUILTIN_MAKE_INTEGER_SEQ(__integer_sequence, size_t, _Ep - _Sp)::template __to_tuple_indices<_Sp>;
 
-#elif defined(_LIBCUDACXX_HAS_INTEGER_PACK) // ^^^ _LIBCUDACXX_HAS_MAKE_INTEGER_SEQ ^^^
-                                            // vvv _LIBCUDACXX_HAS_INTEGER_PACK vvv
+#elif defined(_CCCL_BUILTIN_INTEGER_PACK)
 
 template <size_t _Ep, size_t _Sp>
 using __make_indices_imp _LIBCUDACXX_NODEBUG_TYPE =
-  typename __integer_sequence<size_t, __integer_pack(_Ep - _Sp)...>::template __to_tuple_indices<_Sp>;
+  typename __integer_sequence<size_t, _CCCL_BUILTIN_INTEGER_PACK(_Ep - _Sp)...>::template __to_tuple_indices<_Sp>;
 
-#else // ^^^ _LIBCUDACXX_HAS_INTEGER_PACK ^^^ / vvv !_LIBCUDACXX_HAS_INTEGER_PACK vvv
+#else // ^^^ _CCCL_BUILTIN_INTEGER_PACK ^^^ / vvv !_CCCL_BUILTIN_INTEGER_PACK vvv
 
 namespace __detail
 {
@@ -188,7 +187,7 @@ template <size_t _Ep, size_t _Sp>
 using __make_indices_imp _LIBCUDACXX_NODEBUG_TYPE =
   typename __detail::__make<_Ep - _Sp>::type::template __to_tuple_indices<_Sp>;
 
-#endif
+#endif // !_CCCL_BUILTIN_INTEGER_PACK
 
 template <class _Tp, _Tp... _Ip>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT integer_sequence
@@ -204,18 +203,17 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT integer_sequence
 template <size_t... _Ip>
 using index_sequence = integer_sequence<size_t, _Ip...>;
 
-#ifdef _LIBCUDACXX_HAS_MAKE_INTEGER_SEQ
+#if defined(_CCCL_BUILTIN_MAKE_INTEGER_SEQ)
 
 template <class _Tp, _Tp _Ep>
-using __make_integer_sequence _LIBCUDACXX_NODEBUG_TYPE = __make_integer_seq<integer_sequence, _Tp, _Ep>;
+using __make_integer_sequence _LIBCUDACXX_NODEBUG_TYPE = _CCCL_BUILTIN_MAKE_INTEGER_SEQ(integer_sequence, _Tp, _Ep);
 
-#elif defined(_LIBCUDACXX_HAS_INTEGER_PACK) // ^^^ _LIBCUDACXX_HAS_MAKE_INTEGER_SEQ ^^^
-                                            // vvv _LIBCUDACXX_HAS_INTEGER_PACK vvv
+#elif defined(_CCCL_BUILTIN_INTEGER_PACK)
 
 template <class _Tp, _Tp _Ep>
 using __make_integer_sequence _LIBCUDACXX_NODEBUG_TYPE = integer_sequence<_Tp, __integer_pack(_Ep)...>;
 
-#else // ^^^ _LIBCUDACXX_HAS_INTEGER_PACK ^^^ / vvv !_LIBCUDACXX_HAS_INTEGER_PACK vvv
+#else // ^^^ _CCCL_BUILTIN_INTEGER_PACK ^^^ / vvv !_CCCL_BUILTIN_INTEGER_PACK vvv
 
 template <typename _Tp, _Tp _Np>
 using __make_integer_sequence_unchecked _LIBCUDACXX_NODEBUG_TYPE =
@@ -234,7 +232,7 @@ struct __make_integer_sequence_checked
 template <class _Tp, _Tp _Ep>
 using __make_integer_sequence _LIBCUDACXX_NODEBUG_TYPE = typename __make_integer_sequence_checked<_Tp, _Ep>::type;
 
-#endif // _LIBCUDACXX_HAS_MAKE_INTEGER_SEQ
+#endif // !_CCCL_BUILTIN_INTEGER_PACK
 
 template <class _Tp, _Tp _Np>
 using make_integer_sequence = __make_integer_sequence<_Tp, _Np>;
diff --git a/libcudacxx/include/cuda/std/detail/__annotated_ptr b/libcudacxx/include/cuda/std/detail/__annotated_ptr
index 2b29905af0..3ba339a254 100644
--- a/libcudacxx/include/cuda/std/detail/__annotated_ptr
+++ b/libcudacxx/include/cuda/std/detail/__annotated_ptr
@@ -141,11 +141,11 @@ _CCCL_DEVICE void* __associate_address_space(void* __ptr, _Property __prop)
   {
     bool __b = __isShared(__ptr);
     _LIBCUDACXX_ASSERT(__b, "");
-#if !defined(_CCCL_CUDACC_BELOW_11_2)
-    __builtin_assume(__b);
-#else // ^^^ !_CCCL_CUDACC_BELOW_11_2 ^^^ / vvv _CCCL_CUDACC_BELOW_11_2 vvv
+#if defined(_CCCL_BUILTIN_ASSUME)
+    _CCCL_BUILTIN_ASSUME(__b);
+#else // ^^^ _CCCL_BUILTIN_ASSUME ^^^ / vvv !_CCCL_BUILTIN_ASSUME vvv
     (void) __b;
-#endif // _CCCL_CUDACC_BELOW_11_2
+#endif // !_CCCL_BUILTIN_ASSUME
   }
   else if (std::is_same<_Property, access_property::global>::value == true
            || std::is_same<_Property, access_property::normal>::value == true
@@ -155,11 +155,11 @@ _CCCL_DEVICE void* __associate_address_space(void* __ptr, _Property __prop)
   {
     bool __b = __isGlobal(__ptr);
     _LIBCUDACXX_ASSERT(__b, "");
-#if !defined(_CCCL_CUDACC_BELOW_11_2)
-    __builtin_assume(__b);
-#else // ^^^ !_CCCL_CUDACC_BELOW_11_2 ^^^ / vvv _CCCL_CUDACC_BELOW_11_2 vvv
+#if defined(_CCCL_BUILTIN_ASSUME)
+    _CCCL_BUILTIN_ASSUME(__b);
+#else // ^^^ _CCCL_BUILTIN_ASSUME ^^^ / vvv !_CCCL_BUILTIN_ASSUME vvv
     (void) __b;
-#endif // _CCCL_CUDACC_BELOW_11_2
+#endif // !_CCCL_BUILTIN_ASSUME
   }
 
   return __ptr;
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__assert b/libcudacxx/include/cuda/std/detail/libcxx/include/__assert
index 3568b3b746..61c752cea1 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/__assert
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__assert
@@ -55,9 +55,9 @@
       (void)0 :                                                                 \
       ::_CUDA_VSTD::__libcpp_verbose_abort("%s:%d: assertion %s failed: %s", __FILE__, __LINE__, #expression, message)
     _CCCL_DIAG_POP)
-#elif 0 // !defined(_LIBCUDACXX_ASSERTIONS_DISABLE_ASSUME) && __has_builtin(__builtin_assume)
-#  define _LIBCUDACXX_ASSERT(expression, message)                                                          \
-    (_CCCL_DIAG_PUSH _CCCL_DIAG_SUPPRESS_CLANG("-Wassume") __builtin_assume(static_cast<bool>(expression)) \
+#elif 0 // !defined(_LIBCUDACXX_ASSERTIONS_DISABLE_ASSUME) && defined(_CCCL_BUILTIN_ASSUME)
+#  define _LIBCUDACXX_ASSERT(expression, message)                                                              \
+    (_CCCL_DIAG_PUSH _CCCL_DIAG_SUPPRESS_CLANG("-Wassume") _CCCL_BUILTIN_ASSUME(static_cast<bool>(expression)) \
        _CCCL_DIAG_POP)
 #else
 #  define _LIBCUDACXX_ASSERT(expression, message) ((void) 0)
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__config b/libcudacxx/include/cuda/std/detail/libcxx/include/__config
index bd41943c7e..234dca01f2 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/__config
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__config
@@ -450,13 +450,6 @@ typedef __char32_t char32_t;
 
 #  elif defined(_CCCL_COMPILER_GCC)
 
-#    ifndef _LIBCUDACXX_USE_IS_ASSIGNABLE_FALLBACK
-// FIXME: GCC 8.0 supports this trait, but it has a bug.
-// https://gcc.gnu.org/bugzilla/show_bug.cgi?id=91592
-// https://godbolt.org/z/IljfIw
-#      define _LIBCUDACXX_USE_IS_ASSIGNABLE_FALLBACK
-#    endif // _LIBCUDACXX_USE_IS_ASSIGNABLE_FALLBACK
-
 // GCC 5 supports variable templates
 #    if !defined(__cpp_variable_templates) || __cpp_variable_templates < 201304L
 #      define _LIBCUDACXX_HAS_NO_VARIABLE_TEMPLATES
@@ -705,23 +698,6 @@ typedef unsigned int char32_t;
 #    define _LIBCUDACXX_EXPLICIT
 #  endif
 
-#  if !__has_builtin(__builtin_operator_new) || !__has_builtin(__builtin_operator_delete)
-#    define _LIBCUDACXX_HAS_NO_BUILTIN_OPERATOR_NEW_DELETE
-#  endif
-
-#  if !defined(_LIBCUDACXX_TESTING_FALLBACK_MAKE_INTEGER_SEQUENCE)
-#    if __check_builtin(make_integer_seq)
-#      define _LIBCUDACXX_HAS_MAKE_INTEGER_SEQ
-#    elif defined(_CCCL_COMPILER_MSVC)
-#      if _CCCL_MSVC_VERSION_FULL >= 190023918
-#        define _LIBCUDACXX_HAS_MAKE_INTEGER_SEQ
-#      endif
-#    endif
-#    if __check_builtin(integer_pack)
-#      define _LIBCUDACXX_HAS_INTEGER_PACK
-#    endif
-#  endif
-
 #  ifdef _LIBCUDACXX_DEBUG
 #    if _LIBCUDACXX_DEBUG == 0
 #      define _LIBCUDACXX_DEBUG_LEVEL 1
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__string b/libcudacxx/include/cuda/std/detail/libcxx/include/__string
index 1d6d75a71a..eed3364292 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/__string
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__string
@@ -139,7 +139,8 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT char_traits
 };
 
 template <class _CharT>
-_CCCL_CONSTEXPR_CXX14 int char_traits<_CharT>::compare(const char_type* __s1, const char_type* __s2, size_t __n)
+_LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 int
+char_traits<_CharT>::compare(const char_type* __s1, const char_type* __s2, size_t __n)
 {
   for (; __n; --__n, ++__s1, ++__s2)
   {
@@ -156,7 +157,7 @@ _CCCL_CONSTEXPR_CXX14 int char_traits<_CharT>::compare(const char_type* __s1, co
 }
 
 template <class _CharT>
-inline _CCCL_CONSTEXPR_CXX14 size_t char_traits<_CharT>::length(const char_type* __s)
+_LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 size_t char_traits<_CharT>::length(const char_type* __s)
 {
   size_t __len = 0;
   for (; !eq(*__s, char_type(0)); ++__s)
@@ -167,7 +168,7 @@ inline _CCCL_CONSTEXPR_CXX14 size_t char_traits<_CharT>::length(const char_type*
 }
 
 template <class _CharT>
-inline _CCCL_CONSTEXPR_CXX14 const _CharT*
+_LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 const _CharT*
 char_traits<_CharT>::find(const char_type* __s, size_t __n, const char_type& __a)
 {
   for (; __n; --__n)
@@ -324,7 +325,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT char_traits<char>
 #endif // !__cuda_std__
 };
 
-inline _CCCL_CONSTEXPR_CXX14 int
+_LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 int
 char_traits<char>::compare(const char_type* __s1, const char_type* __s2, size_t __n) noexcept
 {
   if (__n == 0)
@@ -349,7 +350,7 @@ char_traits<char>::compare(const char_type* __s1, const char_type* __s2, size_t
 #endif // !has_feature(constexpr_string_builtins)
 }
 
-inline _CCCL_CONSTEXPR_CXX14 const char*
+_LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 const char*
 char_traits<char>::find(const char_type* __s, size_t __n, const char_type& __a) noexcept
 {
   if (__n == 0)
@@ -702,7 +703,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT char_traits<char16_t>
   }
 };
 
-inline _CCCL_CONSTEXPR_CXX14 int
+_LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 int
 char_traits<char16_t>::compare(const char_type* __s1, const char_type* __s2, size_t __n) noexcept
 {
   for (; __n; --__n, ++__s1, ++__s2)
@@ -719,7 +720,7 @@ char_traits<char16_t>::compare(const char_type* __s1, const char_type* __s2, siz
   return 0;
 }
 
-inline _CCCL_CONSTEXPR_CXX14 size_t char_traits<char16_t>::length(const char_type* __s) noexcept
+_LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 size_t char_traits<char16_t>::length(const char_type* __s) noexcept
 {
   size_t __len = 0;
   for (; !eq(*__s, char_type(0)); ++__s)
@@ -729,7 +730,7 @@ inline _CCCL_CONSTEXPR_CXX14 size_t char_traits<char16_t>::length(const char_typ
   return __len;
 }
 
-inline _CCCL_CONSTEXPR_CXX14 const char16_t*
+_LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 const char16_t*
 char_traits<char16_t>::find(const char_type* __s, size_t __n, const char_type& __a) noexcept
 {
   for (; __n; --__n)
@@ -843,7 +844,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT char_traits<char32_t>
   }
 };
 
-inline _CCCL_CONSTEXPR_CXX14 int
+_LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 int
 char_traits<char32_t>::compare(const char_type* __s1, const char_type* __s2, size_t __n) noexcept
 {
   for (; __n; --__n, ++__s1, ++__s2)
@@ -860,7 +861,7 @@ char_traits<char32_t>::compare(const char_type* __s1, const char_type* __s2, siz
   return 0;
 }
 
-inline _CCCL_CONSTEXPR_CXX14 size_t char_traits<char32_t>::length(const char_type* __s) noexcept
+_LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 size_t char_traits<char32_t>::length(const char_type* __s) noexcept
 {
   size_t __len = 0;
   for (; !eq(*__s, char_type(0)); ++__s)
@@ -870,7 +871,7 @@ inline _CCCL_CONSTEXPR_CXX14 size_t char_traits<char32_t>::length(const char_typ
   return __len;
 }
 
-inline _CCCL_CONSTEXPR_CXX14 const char32_t*
+_LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 const char32_t*
 char_traits<char32_t>::find(const char_type* __s, size_t __n, const char_type& __a) noexcept
 {
   for (; __n; --__n)
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__verbose_abort b/libcudacxx/include/cuda/std/detail/libcxx/include/__verbose_abort
index b44c0506ef..bc20bff968 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/__verbose_abort
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__verbose_abort
@@ -40,7 +40,7 @@ _CCCL_NORETURN _LIBCUDACXX_ATTRIBUTE_FORMAT(__printf__, 1, 2)
   _LIBCUDACXX_HIDE_FROM_ABI void __libcpp_verbose_abort(const char*, ...)
 {
   ::abort();
-  __builtin_unreachable(); // never reached, but needed to tell the compiler that the function never returns
+  _CCCL_UNREACHABLE(); // never reached, but needed to tell the compiler that the function never returns
 }
 
 _LIBCUDACXX_END_NAMESPACE_STD
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/cmath b/libcudacxx/include/cuda/std/detail/libcxx/include/cmath
index 70a064a0fb..f2116e40a1 100644
--- a/libcudacxx/include/cuda/std/detail/libcxx/include/cmath
+++ b/libcudacxx/include/cuda/std/detail/libcxx/include/cmath
@@ -626,12 +626,12 @@ __constexpr_isnan(_A1 __lcpp_x) noexcept
 {
 #if defined(_CCCL_CUDACC_BELOW_11_8)
   return __isnan(__lcpp_x);
-#elif __has_builtin(__builtin_isnan)
+#elif defined(_CCCL_BUILTIN_ISNAN)
   // nvcc at times has issues determining the type of __lcpp_x
-  return __builtin_isnan(static_cast<double>(__lcpp_x));
-#else
+  return _CCCL_BUILTIN_ISNAN(static_cast<double>(__lcpp_x));
+#else // ^^^ _CCCL_BUILTIN_ISNAN ^^^ / vvv !_CCCL_BUILTIN_ISNAN vvv
   return ::isnan(__lcpp_x);
-#endif
+#endif // !_CCCL_BUILTIN_ISNAN
 }
 
 template <class _A1>
@@ -647,12 +647,12 @@ __constexpr_isinf(_A1 __lcpp_x) noexcept
 {
 #if defined(_CCCL_CUDACC_BELOW_11_8)
   return __isinf(__lcpp_x);
-#elif __has_builtin(__builtin_isinf)
+#elif defined(_CCCL_BUILTIN_ISINF)
   // nvcc at times has issues determining the type of __lcpp_x
   return __builtin_isinf(static_cast<double>(__lcpp_x));
-#else
+#else // ^^^ _CCCL_BUILTIN_ISINF ^^^ / vvv !_CCCL_BUILTIN_ISINF vvv
   return ::isinf(__lcpp_x);
-#endif
+#endif // !_CCCL_BUILTIN_ISINF
 }
 
 template <class _A1>
@@ -668,12 +668,12 @@ __constexpr_isfinite(_A1 __lcpp_x) noexcept
 {
 #if defined(_CCCL_CUDACC_BELOW_11_8)
   return !__isinf(__lcpp_x) && !__isnan(__lcpp_x);
-#elif __has_builtin(__builtin_isfinite)
+#elif defined(_CCCL_BUILTIN_ISFINITE)
   // nvcc at times has issues determining the type of __lcpp_x
   return __builtin_isfinite(static_cast<double>(__lcpp_x));
-#else
+#else // ^^^ _CCCL_BUILTIN_ISFINITE ^^^ / vvv !_CCCL_BUILTIN_ISFINITE vvv
   return ::isfinite(__lcpp_x);
-#endif
+#endif // !_CCCL_BUILTIN_ISFINITE
 }
 
 template <class _A1>
diff --git a/libcudacxx/test/libcudacxx/std/utilities/function.objects/func.not_fn/not_fn.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/function.objects/func.not_fn/not_fn.pass.cpp
index 943130c41d..543914731e 100644
--- a/libcudacxx/test/libcudacxx/std/utilities/function.objects/func.not_fn/not_fn.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/utilities/function.objects/func.not_fn/not_fn.pass.cpp
@@ -505,11 +505,7 @@ void throws_in_constructor_test()
     bool operator()() const
     {
       assert(false);
-#  if defined(TEST_COMPILER_MSVC)
-      __assume(0);
-#  else
-      __builtin_unreachable();
-#  endif
+      _CCCL_UNREACHABLE();
     }
   };
   {
diff --git a/libcudacxx/test/libcudacxx/std/utilities/meta/meta.trans/meta.trans.ptr/add_pointer.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/meta/meta.trans/meta.trans.ptr/add_pointer.pass.cpp
index 586ebfc609..0f36c87993 100644
--- a/libcudacxx/test/libcudacxx/std/utilities/meta/meta.trans/meta.trans.ptr/add_pointer.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/utilities/meta/meta.trans/meta.trans.ptr/add_pointer.pass.cpp
@@ -23,27 +23,27 @@ template <class T, class U>
 __host__ __device__ void test_add_pointer()
 {
   ASSERT_SAME_TYPE(U, typename cuda::std::add_pointer<T>::type);
-#if TEST_STD_VER > 2011
+#if TEST_STD_VER >= 2014
   ASSERT_SAME_TYPE(U, cuda::std::add_pointer_t<T>);
-#endif
+#endif // TEST_STD_VER >= 2014
 }
 
 template <class F>
 __host__ __device__ void test_function0()
 {
   ASSERT_SAME_TYPE(F*, typename cuda::std::add_pointer<F>::type);
-#if TEST_STD_VER > 2011
+#if TEST_STD_VER >= 2014
   ASSERT_SAME_TYPE(F*, cuda::std::add_pointer_t<F>);
-#endif
+#endif // TEST_STD_VER >= 2014
 }
 
 template <class F>
 __host__ __device__ void test_function1()
 {
   ASSERT_SAME_TYPE(F, typename cuda::std::add_pointer<F>::type);
-#if TEST_STD_VER > 2011
+#if TEST_STD_VER >= 2014
   ASSERT_SAME_TYPE(F, cuda::std::add_pointer_t<F>);
-#endif
+#endif // TEST_STD_VER >= 2014
 }
 
 struct Foo
diff --git a/libcudacxx/test/libcudacxx/std/utilities/optional/optional.monadic/or_else.pass.cpp b/libcudacxx/test/libcudacxx/std/utilities/optional/optional.monadic/or_else.pass.cpp
index 57732aa7f8..9f915526b2 100644
--- a/libcudacxx/test/libcudacxx/std/utilities/optional/optional.monadic/or_else.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/utilities/optional/optional.monadic/or_else.pass.cpp
@@ -95,10 +95,7 @@ __host__ __device__ TEST_CONSTEXPR_CXX17 bool test()
     opt = 1;
     opt.or_else([] {
 #if defined(TEST_COMPILER_GCC) && __GNUC__ < 9
-      // This is an awful hack around: `error: call to non-'constexpr' function 'void __assert_fail(...)'`
-      // Verified experimentally that __builtin_unreachable is treated the way we want, i.e. compiles
-      // fine within the context of a constexpr expression if unevaluated, but generates an error when evaluated.
-      NV_IF_TARGET(NV_IS_HOST, (__builtin_unreachable();), (__trap();))
+      _CCCL_UNREACHABLE();
 #else
       assert(false);
 #endif