Expose more of construct at earlier

docs/libcudacxx/standard_api/utility_library/memory.rst

@@ -17,10 +17,11 @@ Extensions
 ----------
 
 -  Most features are available from C++11 onwards.
--  All features are available at compile time if compiler support is sufficient.
+-  ``cuda::std::addressof`` is constexpr from C++11 on if compiler support is available
+-  ``cuda::std::assume_aligned`` is constexpr from C++14 on
 
 Restrictions
 ------------
 
-- The features that are explicitly named in the standard `construct_at` and `destroy_at` are only available in C++20
+- `construct_at` and is only available in C++20 as that is explicitly mentioned in the standard
 - The specialized memory algorithms are not parallel

libcudacxx/include/cuda/std/__memory/assume_aligned.h

@@ -30,7 +30,7 @@
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
 template <size_t _Align, class _Tp>
-_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp* assume_aligned(_Tp* __ptr) noexcept
+_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 _Tp* assume_aligned(_Tp* __ptr) noexcept
 {
   static_assert(_CUDA_VSTD::has_single_bit(_Align), "std::assume_aligned requires the alignment to be a power of 2!");
   if (_CUDA_VSTD::__libcpp_is_constant_evaluated())

libcudacxx/include/cuda/std/__memory/construct_at.h

@@ -4,7 +4,7 @@
 // Part of the LLVM Project, 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) 2023 NVIDIA CORPORATION & AFFILIATES
+// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES
 //
 //===----------------------------------------------------------------------===//
 
@@ -185,8 +185,8 @@ _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 _ForwardIterator __destroy(_Forw
 
 _CCCL_EXEC_CHECK_DISABLE
 template <class _Tp,
-          __enable_if_t<!is_array<_Tp>::value, int>                  = 0,
-          __enable_if_t<!is_trivially_destructible<_Tp>::value, int> = 0>
+          __enable_if_t<!_CCCL_TRAIT(is_array, _Tp), int>                  = 0,
+          __enable_if_t<!_CCCL_TRAIT(is_trivially_destructible, _Tp), int> = 0>
 _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 void __destroy_at(_Tp* __loc)
 {
   _CCCL_ASSERT(__loc != nullptr, "null pointer given to destroy_at");
@@ -195,22 +195,20 @@ _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 void __destroy_at(_Tp* __loc)
 
 _CCCL_EXEC_CHECK_DISABLE
 template <class _Tp,
-          __enable_if_t<!is_array<_Tp>::value, int>                 = 0,
-          __enable_if_t<is_trivially_destructible<_Tp>::value, int> = 0>
+          __enable_if_t<!_CCCL_TRAIT(is_array, _Tp), int>                 = 0,
+          __enable_if_t<_CCCL_TRAIT(is_trivially_destructible, _Tp), int> = 0>
 _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 void __destroy_at(_Tp* __loc)
 {
   _CCCL_ASSERT(__loc != nullptr, "null pointer given to destroy_at");
   (void) __loc;
 }
 
-#if _CCCL_STD_VER >= 2020
-template <class _Tp, __enable_if_t<is_array<_Tp>::value, int> = 0>
-_LIBCUDACXX_HIDE_FROM_ABI constexpr void __destroy_at(_Tp* __loc)
+template <class _Tp, __enable_if_t<_CCCL_TRAIT(is_array, _Tp), int> = 0>
+_LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 void __destroy_at(_Tp* __loc)
 {
   _CCCL_ASSERT(__loc != nullptr, "null pointer given to destroy_at");
   _CUDA_VSTD::__destroy(_CUDA_VSTD::begin(*__loc), _CUDA_VSTD::end(*__loc));
 }
-#endif // _CCCL_STD_VER >= 2020
 
 template <class _ForwardIterator>
 _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 _ForwardIterator
@@ -235,22 +233,18 @@ __reverse_destroy(_BidirectionalIterator __first, _BidirectionalIterator __last)
   return __last;
 }
 
-#if _CCCL_STD_VER >= 2017
-
-template <class _Tp, enable_if_t<!is_array_v<_Tp>, int> = 0>
-_LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX20 void destroy_at(_Tp* __loc) noexcept
+template <class _Tp, __enable_if_t<!_CCCL_TRAIT(is_array, _Tp), int> = 0>
+_LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX20 void destroy_at(_Tp* __loc)
 {
   _CCCL_ASSERT(__loc != nullptr, "null pointer given to destroy_at");
   __loc->~_Tp();
 }
 
-#  if _CCCL_STD_VER >= 2020
-template <class _Tp, enable_if_t<is_array_v<_Tp>, int> = 0>
-_LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX20 void destroy_at(_Tp* __loc) noexcept
+template <class _Tp, __enable_if_t<_CCCL_TRAIT(is_array, _Tp), int> = 0>
+_LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX20 void destroy_at(_Tp* __loc)
 {
   _CUDA_VSTD::__destroy_at(__loc);
 }
-#  endif // _CCCL_STD_VER >= 2020
 
 template <class _ForwardIterator>
 _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX20 void destroy(_ForwardIterator __first, _ForwardIterator __last) noexcept
@@ -268,8 +262,6 @@ _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX20 _ForwardIterator destroy_n(_Forw
   return __first;
 }
 
-#endif // _CCCL_STD_VER >= 2017
-
 _LIBCUDACXX_END_NAMESPACE_STD
 
 #endif // _LIBCUDACXX___MEMORY_CONSTRUCT_AT_H

libcudacxx/test/libcudacxx/std/utilities/memory/ptr.align/assume_aligned.pass.cpp

@@ -19,40 +19,40 @@
 #include "test_macros.h"
 
 template <typename T>
-__host__ __device__ constexpr void check(T* p)
+__host__ __device__ TEST_CONSTEXPR_CXX14 void check(T* p)
 {
   ASSERT_SAME_TYPE(T*, decltype(cuda::std::assume_aligned<1>(p)));
   constexpr cuda::std::size_t alignment = alignof(T);
 
-  if constexpr (alignment >= 1)
+  _CCCL_IF_CONSTEXPR (alignment >= 1)
   {
     assert(p == cuda::std::assume_aligned<1>(p));
   }
-  if constexpr (alignment >= 2)
+  _CCCL_IF_CONSTEXPR (alignment >= 2)
   {
     assert(p == cuda::std::assume_aligned<2>(p));
   }
-  if constexpr (alignment >= 4)
+  _CCCL_IF_CONSTEXPR (alignment >= 4)
   {
     assert(p == cuda::std::assume_aligned<4>(p));
   }
-  if constexpr (alignment >= 8)
+  _CCCL_IF_CONSTEXPR (alignment >= 8)
   {
     assert(p == cuda::std::assume_aligned<8>(p));
   }
-  if constexpr (alignment >= 16)
+  _CCCL_IF_CONSTEXPR (alignment >= 16)
   {
     assert(p == cuda::std::assume_aligned<16>(p));
   }
-  if constexpr (alignment >= 32)
+  _CCCL_IF_CONSTEXPR (alignment >= 32)
   {
     assert(p == cuda::std::assume_aligned<32>(p));
   }
-  if constexpr (alignment >= 64)
+  _CCCL_IF_CONSTEXPR (alignment >= 64)
   {
     assert(p == cuda::std::assume_aligned<64>(p));
   }
-  if constexpr (alignment >= 128)
+  _CCCL_IF_CONSTEXPR (alignment >= 128)
   {
     assert(p == cuda::std::assume_aligned<128>(p));
   }
@@ -73,7 +73,7 @@ struct alignas(64) S64
 struct alignas(128) S128
 {};
 
-__host__ __device__ constexpr bool tests()
+__host__ __device__ TEST_CONSTEXPR_CXX14 bool tests()
 {
   char c{};
   int i{};
@@ -105,7 +105,9 @@ __host__ __device__ constexpr bool tests()
 int main(int, char**)
 {
   tests();
+#if TEST_STD_VER >= 2014
   static_assert(tests());
+#endif // TEST_STD_VER >= 2014
 
   return 0;
 }

libcudacxx/test/libcudacxx/std/utilities/memory/specialized.algorithms/specialized.construct/construct_at.pass.cpp

@@ -16,7 +16,7 @@
 
 // #include <cuda/std/memory>
 #include <cuda/std/cassert>
-#include <cuda/std/utility>
+#include <cuda/std/memory>
 
 #include "test_iterators.h"
 #include "test_macros.h"

libcudacxx/test/libcudacxx/std/utilities/memory/specialized.algorithms/specialized.destroy/destroy.pass.cpp

@@ -7,7 +7,6 @@
 //
 //===----------------------------------------------------------------------===//
 
-// UNSUPPORTED: c++03, c++11, c++14
 // UNSUPPORTED: gcc-6
 
 // <memory>
@@ -17,21 +16,21 @@
 
 // #include <cuda/std/memory>
 #include <cuda/std/cassert>
+#include <cuda/std/memory>
 #include <cuda/std/type_traits>
-#include <cuda/std/utility>
 
 #include "test_iterators.h"
 #include "test_macros.h"
 
 struct Counted
 {
   int* counter_ = nullptr;
-  __host__ __device__ TEST_CONSTEXPR Counted(int* counter)
+  __host__ __device__ TEST_CONSTEXPR_CXX14 Counted(int* counter)
       : counter_(counter)
   {
     ++*counter_;
   }
-  __host__ __device__ TEST_CONSTEXPR Counted(Counted const& other)
+  __host__ __device__ TEST_CONSTEXPR_CXX14 Counted(Counted const& other)
       : counter_(other.counter_)
   {
     ++*counter_;
@@ -43,8 +42,7 @@ struct Counted
   __host__ __device__ friend void operator&(Counted) = delete;
 };
 
-#if TEST_STD_VER > 2017
-__host__ __device__ constexpr bool test_arrays()
+__host__ __device__ TEST_CONSTEXPR_CXX20 bool test_arrays()
 {
   {
     int counter     = 0;
@@ -83,7 +81,6 @@ __host__ __device__ constexpr bool test_arrays()
 
   return true;
 }
-#endif
 
 template <class It>
 __host__ __device__ TEST_CONSTEXPR_CXX20 void test()
@@ -113,8 +110,8 @@ __host__ __device__ TEST_CONSTEXPR_CXX20 bool tests()
 int main(int, char**)
 {
   tests();
-#if TEST_STD_VER > 2017
   test_arrays();
+#if TEST_STD_VER > 2017
 #  if !defined(TEST_COMPILER_NVRTC)
 #    if (defined(TEST_COMPILER_CLANG) && __clang_major__ > 10) || (defined(TEST_COMPILER_GCC) && __GNUC__ > 9) \
       || defined(TEST_COMPILER_MSVC_2022) || defined(TEST_COMPILER_NVHPC)

libcudacxx/test/libcudacxx/std/utilities/memory/specialized.algorithms/specialized.destroy/destroy_at.pass.cpp

@@ -7,7 +7,6 @@
 //
 //===----------------------------------------------------------------------===//
 
-// UNSUPPORTED: c++03, c++11, c++14
 // UNSUPPORTED: gcc-6
 
 // <memory>
@@ -17,15 +16,15 @@
 
 // #include <cuda/std/memory>
 #include <cuda/std/cassert>
+#include <cuda/std/memory>
 #include <cuda/std/type_traits>
-#include <cuda/std/utility>
 
 #include "test_macros.h"
 
 struct Counted
 {
   int* counter_;
-  __host__ __device__ TEST_CONSTEXPR Counted(int* counter)
+  __host__ __device__ TEST_CONSTEXPR_CXX14 Counted(int* counter)
       : counter_(counter)
   {
     ++*counter_;
@@ -40,7 +39,7 @@ struct Counted
 struct VirtualCounted
 {
   int* counter_;
-  __host__ __device__ TEST_CONSTEXPR VirtualCounted(int* counter)
+  __host__ __device__ TEST_CONSTEXPR_CXX14 VirtualCounted(int* counter)
       : counter_(counter)
   {
     ++*counter_;
@@ -54,14 +53,13 @@ struct VirtualCounted
 
 struct DerivedCounted : VirtualCounted
 {
-  __host__ __device__ TEST_CONSTEXPR DerivedCounted(int* counter)
+  __host__ __device__ TEST_CONSTEXPR_CXX14 DerivedCounted(int* counter)
       : VirtualCounted(counter)
   {}
   __host__ __device__ TEST_CONSTEXPR_CXX20 ~DerivedCounted() override {}
 };
 
-#if TEST_STD_VER > 2017
-__host__ __device__ constexpr bool test_arrays()
+__host__ __device__ TEST_CONSTEXPR_CXX20 bool test_arrays()
 {
   {
     int counter    = 0;
@@ -98,7 +96,6 @@ __host__ __device__ constexpr bool test_arrays()
   }
   return true;
 }
-#endif
 
 __host__ __device__ TEST_CONSTEXPR_CXX20 bool test()
 {
@@ -143,8 +140,8 @@ __host__ __device__ TEST_CONSTEXPR_CXX20 bool test()
 int main(int, char**)
 {
   test();
-#if TEST_STD_VER > 2017
   test_arrays();
+#if TEST_STD_VER > 2017
 #  if !defined(TEST_COMPILER_NVRTC)
 #    if (defined(TEST_COMPILER_CLANG) && __clang_major__ > 10) || (defined(TEST_COMPILER_GCC) && __GNUC__ > 9) \
       || defined(TEST_COMPILER_MSVC_2022) || defined(TEST_COMPILER_NVHPC)

libcudacxx/test/libcudacxx/std/utilities/memory/specialized.algorithms/specialized.destroy/destroy_n.pass.cpp

@@ -7,7 +7,6 @@
 //
 //===----------------------------------------------------------------------===//
 
-// UNSUPPORTED: c++03, c++11, c++14
 // UNSUPPORTED: gcc-6
 
 // <memory>
@@ -17,21 +16,21 @@
 
 // #include <cuda/std/memory>
 #include <cuda/std/cassert>
+#include <cuda/std/memory>
 #include <cuda/std/type_traits>
-#include <cuda/std/utility>
 
 #include "test_iterators.h"
 #include "test_macros.h"
 
 struct Counted
 {
   int* counter_ = nullptr;
-  __host__ __device__ TEST_CONSTEXPR Counted(int* counter)
+  __host__ __device__ TEST_CONSTEXPR_CXX14 Counted(int* counter)
       : counter_(counter)
   {
     ++*counter_;
   }
-  __host__ __device__ TEST_CONSTEXPR Counted(Counted const& other)
+  __host__ __device__ TEST_CONSTEXPR_CXX14 Counted(Counted const& other)
       : counter_(other.counter_)
   {
     ++*counter_;
@@ -43,8 +42,7 @@ struct Counted
   __host__ __device__ friend void operator&(Counted) = delete;
 };
 
-#if TEST_STD_VER > 2017
-__host__ __device__ constexpr bool test_arrays()
+__host__ __device__ TEST_CONSTEXPR_CXX20 bool test_arrays()
 {
   {
     int counter     = 0;
@@ -70,7 +68,6 @@ __host__ __device__ constexpr bool test_arrays()
 
   return true;
 }
-#endif
 
 template <class It>
 __host__ __device__ TEST_CONSTEXPR_CXX20 void test()
@@ -102,8 +99,8 @@ __host__ __device__ TEST_CONSTEXPR_CXX20 bool tests()
 int main(int, char**)
 {
   tests();
-#if TEST_STD_VER > 2017
   test_arrays();
+#if TEST_STD_VER > 2017
 #  if !defined(TEST_COMPILER_NVRTC)
 #    if (defined(TEST_COMPILER_CLANG) && __clang_major__ > 10) || (defined(TEST_COMPILER_GCC) && __GNUC__ > 9) \
       || defined(TEST_COMPILER_MSVC_2022) || defined(TEST_COMPILER_NVHPC)