Skip to content

Commit e048b61

Browse files
davebayerdavebayer
authored
Add if consteval support in libcu++ (#6424)
1 parent 16cc647 commit e048b61

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

48 files changed

+241
-171
lines changed

‎cudax/include/cuda/experimental/__launch/configuration.cuh‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -296,7 +296,7 @@ public:
296296
{
297297
if constexpr (::cuda::std::is_unbounded_array_v<_Tp>)
298298
{
299-
if (!::cuda::std::__cccl_default_is_constant_evaluated())
299+
_CCCL_IF_NOT_CONSTEVAL_DEFAULT
300300
{
301301
NV_IF_TARGET(NV_IS_DEVICE, (return ::cuda::ptx::get_sreg_dynamic_smem_size();))
302302
}

‎libcudacxx/include/cuda/__annotated_ptr/access_property_encoding.h‎

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -26,7 +26,6 @@
2626
#include <cuda/std/__algorithm/clamp.h>
2727
#include <cuda/std/__algorithm/max.h>
2828
#include <cuda/std/__bit/bit_cast.h>
29-
#include <cuda/std/__type_traits/is_constant_evaluated.h>
3029
#include <cuda/std/__utility/to_underlying.h>
3130
#include <cuda/std/cstddef>
3231
#include <cuda/std/cstdint>
@@ -138,7 +137,7 @@ enum class __l2_descriptor_mode_t : uint32_t
138137
[[nodiscard]] _CCCL_API constexpr uint64_t
139138
__l2_interleave(__l2_evict_t __primary, __l2_evict_t __secondary, float __fraction)
140139
{
141-
if (!::cuda::std::__cccl_default_is_constant_evaluated())
140+
_CCCL_IF_NOT_CONSTEVAL_DEFAULT
142141
{
143142
NV_IF_ELSE_TARGET(
144143
NV_PROVIDES_SM_80, (return ::cuda::__createpolicy_fraction(__primary, __secondary, __fraction);), (return 0;))

‎libcudacxx/include/cuda/__annotated_ptr/annotated_ptr.h‎

Lines changed: 2 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -25,7 +25,6 @@
2525
#include <cuda/__annotated_ptr/annotated_ptr_base.h>
2626
#include <cuda/__memcpy_async/memcpy_async.h>
2727
#include <cuda/__memory/address_space.h>
28-
#include <cuda/std/__type_traits/is_constant_evaluated.h>
2928
#include <cuda/std/cstddef>
3029
#include <cuda/std/cstdint>
3130

@@ -74,7 +73,7 @@ class annotated_ptr : private ::cuda::__annotated_ptr_base<_Property>
7473
NV_IF_TARGET(NV_IS_HOST, (_CCCL_ASSERT(!__is_smem, "shared memory pointer is not supported on the host");))
7574
if constexpr (__is_smem)
7675
{
77-
if (!::cuda::std::__cccl_default_is_constant_evaluated())
76+
_CCCL_IF_NOT_CONSTEVAL_DEFAULT
7877
{
7978
NV_IF_TARGET(NV_IS_DEVICE,
8079
(_CCCL_ASSERT(::cuda::device::is_address_from(__p, ::cuda::device::address_space::shared),
@@ -84,7 +83,7 @@ class annotated_ptr : private ::cuda::__annotated_ptr_base<_Property>
8483
else
8584
{
8685
_CCCL_ASSERT(__p != nullptr, "__p must not be null");
87-
if (!::cuda::std::__cccl_default_is_constant_evaluated())
86+
_CCCL_IF_NOT_CONSTEVAL_DEFAULT
8887
{
8988
NV_IF_TARGET(NV_IS_DEVICE,
9089
(_CCCL_ASSERT(::cuda::device::is_address_from(__p, ::cuda::device::address_space::global),

‎libcudacxx/include/cuda/__bit/bit_reverse.h‎

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -21,7 +21,6 @@
2121
# pragma system_header
2222
#endif // no system header
2323

24-
#include <cuda/std/__type_traits/is_constant_evaluated.h>
2524
#include <cuda/std/__type_traits/is_unsigned_integer.h>
2625
#include <cuda/std/cstdint>
2726

@@ -181,7 +180,7 @@ template <typename _Tp>
181180
[[nodiscard]] _CCCL_API constexpr _Tp bit_reverse(_Tp __value) noexcept
182181
{
183182
static_assert(::cuda::std::__cccl_is_cv_unsigned_integer_v<_Tp>, "bit_reverse() requires unsigned integer types");
184-
if (!::cuda::std::__cccl_default_is_constant_evaluated())
183+
_CCCL_IF_NOT_CONSTEVAL_DEFAULT
185184
{
186185
NV_IF_TARGET(NV_IS_DEVICE, (return ::cuda::__bit_reverse_device(__value);))
187186
}

‎libcudacxx/include/cuda/__bit/bitfield.h‎

Lines changed: 2 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -24,7 +24,6 @@
2424
#include <cuda/__bit/bitmask.h>
2525
#include <cuda/std/__limits/numeric_limits.h>
2626
#include <cuda/std/__type_traits/conditional.h>
27-
#include <cuda/std/__type_traits/is_constant_evaluated.h>
2827
#include <cuda/std/__type_traits/is_unsigned_integer.h>
2928
#include <cuda/std/cstdint>
3029
#include <cuda/std/limits>
@@ -76,7 +75,7 @@ bitfield_insert(const _Tp __dest, const _Tp __source, int __start, int __width)
7675
_CCCL_ASSERT(__start + __width <= __digits, "start position + width out of range");
7776
if constexpr (sizeof(_Tp) <= sizeof(uint64_t))
7877
{
79-
if (!::cuda::std::__cccl_default_is_constant_evaluated())
78+
_CCCL_IF_NOT_CONSTEVAL_DEFAULT
8079
{
8180
// clang-format off
8281
NV_DISPATCH_TARGET( // all SM < 70
@@ -102,7 +101,7 @@ template <typename _Tp>
102101
_CCCL_ASSERT(__start + __width <= __digits, "start position + width out of range");
103102
if constexpr (sizeof(_Tp) <= sizeof(uint32_t))
104103
{
105-
if (!::cuda::std::__cccl_default_is_constant_evaluated())
104+
_CCCL_IF_NOT_CONSTEVAL_DEFAULT
106105
{
107106
// clang-format off
108107
NV_DISPATCH_TARGET( // all SM < 70

‎libcudacxx/include/cuda/__bit/bitmask.h‎

Lines changed: 3 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -27,7 +27,6 @@
2727
# include <cuda/__ptx/instructions/shr.h>
2828
#endif // _CCCL_CUDA_COMPILATION()
2929
#include <cuda/std/__type_traits/conditional.h>
30-
#include <cuda/std/__type_traits/is_constant_evaluated.h>
3130
#include <cuda/std/__type_traits/is_unsigned_integer.h>
3231
#include <cuda/std/limits>
3332

@@ -38,7 +37,7 @@ _CCCL_BEGIN_NAMESPACE_CUDA
3837
template <typename _Tp>
3938
[[nodiscard]] _CCCL_API constexpr _Tp __shl(const _Tp __value, int __shift) noexcept
4039
{
41-
if (!::cuda::std::__cccl_default_is_constant_evaluated())
40+
_CCCL_IF_NOT_CONSTEVAL_DEFAULT
4241
{
4342
if constexpr (sizeof(_Tp) <= sizeof(uint64_t))
4443
{
@@ -53,7 +52,7 @@ template <typename _Tp>
5352
template <typename _Tp>
5453
[[nodiscard]] _CCCL_API constexpr _Tp __shr(const _Tp __value, int __shift) noexcept
5554
{
56-
if (!::cuda::std::__cccl_default_is_constant_evaluated())
55+
_CCCL_IF_NOT_CONSTEVAL_DEFAULT
5756
{
5857
if constexpr (sizeof(_Tp) <= sizeof(uint64_t))
5958
{
@@ -73,7 +72,7 @@ template <typename _Tp = uint32_t>
7372
_CCCL_ASSERT(__width >= 0 && __width <= __digits, "width out of range");
7473
_CCCL_ASSERT(__start >= 0 && __start <= __digits, "start position out of range");
7574
_CCCL_ASSERT(__start + __width <= __digits, "start position + width out of range");
76-
if (!::cuda::std::__cccl_default_is_constant_evaluated())
75+
_CCCL_IF_NOT_CONSTEVAL_DEFAULT
7776
{
7877
if constexpr (sizeof(_Tp) <= sizeof(uint32_t))
7978
{

‎libcudacxx/include/cuda/__cmath/ceil_div.h‎

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -24,7 +24,6 @@
2424
#include <cuda/std/__algorithm/min.h>
2525
#include <cuda/std/__concepts/concept_macros.h>
2626
#include <cuda/std/__type_traits/common_type.h>
27-
#include <cuda/std/__type_traits/is_constant_evaluated.h>
2827
#include <cuda/std/__type_traits/is_enum.h>
2928
#include <cuda/std/__type_traits/is_integral.h>
3029
#include <cuda/std/__type_traits/is_signed.h>
@@ -61,7 +60,7 @@ _CCCL_REQUIRES(::cuda::std::is_integral_v<_Tp> _CCCL_AND ::cuda::std::is_integra
6160
}
6261
else
6362
{
64-
if (::cuda::std::is_constant_evaluated())
63+
_CCCL_IF_CONSTEVAL
6564
{
6665
const auto __res = __a1 / __b1;
6766
return static_cast<_Common>(__res + (__res * __b1 != __a1));

‎libcudacxx/include/cuda/__cmath/ilog.h‎

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -26,7 +26,6 @@
2626
#include <cuda/std/__cmath/rounding_functions.h>
2727
#include <cuda/std/__concepts/concept_macros.h>
2828
#include <cuda/std/__limits/numeric_limits.h>
29-
#include <cuda/std/__type_traits/is_constant_evaluated.h>
3029
#include <cuda/std/__type_traits/is_integer.h>
3130
#include <cuda/std/__type_traits/is_same.h>
3231
#include <cuda/std/__type_traits/make_unsigned.h>

‎libcudacxx/include/cuda/__cmath/mul_hi.h‎

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -21,7 +21,6 @@
2121
# pragma system_header
2222
#endif // no system header
2323

24-
#include <cuda/std/__type_traits/is_constant_evaluated.h>
2524
#include <cuda/std/__type_traits/is_integer.h>
2625
#include <cuda/std/__type_traits/is_signed.h>
2726
#include <cuda/std/__type_traits/make_nbit_int.h>
@@ -68,7 +67,7 @@ _CCCL_API constexpr _Tp mul_hi(_Tp __lhs, _Tp __rhs) noexcept
6867
{
6968
using ::cuda::std::int64_t;
7069
using ::cuda::std::is_signed_v;
71-
if (!::cuda::std::__cccl_default_is_constant_evaluated())
70+
_CCCL_IF_NOT_CONSTEVAL_DEFAULT
7271
{
7372
if constexpr (sizeof(_Tp) == sizeof(int))
7473
{

‎libcudacxx/include/cuda/__device/device_ref.h‎

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -26,7 +26,6 @@
2626
# include <cuda/__driver/driver_api.h>
2727
# include <cuda/__fwd/devices.h>
2828
# include <cuda/__runtime/types.h>
29-
# include <cuda/std/__type_traits/is_constant_evaluated.h>
3029
# include <cuda/std/span>
3130
# include <cuda/std/string_view>
3231

@@ -46,7 +45,7 @@ class device_ref
4645
/*implicit*/ _CCCL_HOST_API constexpr device_ref(int __id) noexcept
4746
: __id_(__id)
4847
{
49-
if (::cuda::std::__cccl_default_is_constant_evaluated())
48+
_CCCL_IF_CONSTEVAL_DEFAULT
5049
{
5150
_CCCL_VERIFY(__id >= 0, "Device ID must be a valid GPU device ordinal");
5251
}

0 commit comments

Comments
 (0)