Skip to content

Commit 87f7246

Browse files
misccomiscco
authored
We currently do not expose any meow_t alias until C++14 (NVIDIA#2763)
This forces us to duplicate a ton of code with `__meow_t` for no benbefit at all We already backport half the standard so kets just do the right thing and use those aliases whenever possible While we are at it also actually use the `_CCCL_NODDEBUG_ALIAS` attribute when possible
1 parent 16ac5be commit 87f7246

File tree

245 files changed

+1319
-1497
lines changed

Some content is hidden

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

245 files changed

+1319
-1497
lines changed

‎c2h/include/c2h/catch2_test_helper.h

+2-2
Original file line numberDiff line numberDiff line change
@@ -193,11 +193,11 @@ auto BitwiseEqualsRange(const Range& range) -> CustomEqualsRangeMatcher<Range, b
193193
#include <cuda/std/tuple>
194194
_LIBCUDACXX_BEGIN_NAMESPACE_STD
195195
template <size_t N, typename... T>
196-
__enable_if_t<(N == sizeof...(T))> print_elem(::std::ostream&, const tuple<T...>&)
196+
enable_if_t<(N == sizeof...(T))> print_elem(::std::ostream&, const tuple<T...>&)
197197
{}
198198

199199
template <size_t N, typename... T>
200-
__enable_if_t<(N < sizeof...(T))> print_elem(::std::ostream& os, const tuple<T...>& tup)
200+
enable_if_t<(N < sizeof...(T))> print_elem(::std::ostream& os, const tuple<T...>& tup)
201201
{
202202
_CCCL_IF_CONSTEXPR (N != 0)
203203
{

‎cub/benchmarks/bench/transform/babelstream.h

+1-1
Original file line numberDiff line numberDiff line change
@@ -91,7 +91,7 @@ struct narrowing_error : std::runtime_error
9191

9292
// from C++ GSL
9393
// implementation insipired by: https://github.com/microsoft/GSL/blob/main/include/gsl/narrow
94-
template <typename DstT, typename SrcT, ::cuda::std::__enable_if_t<::cuda::std::is_arithmetic<SrcT>::value, int> = 0>
94+
template <typename DstT, typename SrcT, ::cuda::std::enable_if_t<::cuda::std::is_arithmetic<SrcT>::value, int> = 0>
9595
constexpr DstT narrow(SrcT value)
9696
{
9797
constexpr bool is_different_signedness = ::cuda::std::is_signed<SrcT>::value != ::cuda::std::is_signed<DstT>::value;

‎cub/cub/detail/choose_offset.cuh

+1-1
Original file line numberDiff line numberDiff line change
@@ -148,7 +148,7 @@ using choose_signed_offset_t = typename choose_signed_offset<NumItemsT>::type;
148148
template <typename... Iter>
149149
struct common_iterator_value
150150
{
151-
using type = ::cuda::std::__common_type_t<::cuda::std::__iter_value_type<Iter>...>;
151+
using type = ::cuda::std::common_type_t<::cuda::std::__iter_value_type<Iter>...>;
152152
};
153153
template <typename... Iter>
154154
using common_iterator_value_t = typename common_iterator_value<Iter...>::type;

‎cub/cub/detail/type_traits.cuh

+1-1
Original file line numberDiff line numberDiff line change
@@ -50,7 +50,7 @@ _CCCL_SUPPRESS_DEPRECATED_PUSH
5050
_CCCL_SUPPRESS_DEPRECATED_POP
5151
#include <cuda/std/type_traits>
5252

53-
#define _CUB_TEMPLATE_REQUIRES(...) ::cuda::std::__enable_if_t<(__VA_ARGS__)>* = nullptr
53+
#define _CUB_TEMPLATE_REQUIRES(...) ::cuda::std::enable_if_t<(__VA_ARGS__)>* = nullptr
5454

5555
CUB_NAMESPACE_BEGIN
5656
namespace detail

‎cub/cub/device/dispatch/dispatch_merge.cuh

+1-1
Original file line numberDiff line numberDiff line change
@@ -40,7 +40,7 @@ class choose_merge_agent
4040
&& sizeof(typename fallback_agent_t::TempStorage) <= max_smem_per_block;
4141

4242
public:
43-
using type = ::cuda::std::__conditional_t<use_fallback, fallback_agent_t, default_agent_t>;
43+
using type = ::cuda::std::conditional_t<use_fallback, fallback_agent_t, default_agent_t>;
4444
};
4545

4646
// Computes the merge path intersections at equally wide intervals. The approach is outlined in the paper:

‎cub/cub/device/dispatch/dispatch_transform.cuh

+6-6
Original file line numberDiff line numberDiff line change
@@ -147,7 +147,7 @@ _CCCL_DEVICE _CCCL_FORCEINLINE void prefetch_tile(const T* addr, int tile_size)
147147
// TODO(miscco): we should probably constrain It to not be a contiguous iterator in C++17 (and change the overload
148148
// above to accept any contiguous iterator)
149149
// overload for any iterator that is not a pointer, do nothing
150-
template <int, typename It, ::cuda::std::__enable_if_t<!::cuda::std::is_pointer<It>::value, int> = 0>
150+
template <int, typename It, ::cuda::std::enable_if_t<!::cuda::std::is_pointer<It>::value, int> = 0>
151151
_CCCL_DEVICE _CCCL_FORCEINLINE void prefetch_tile(It, int)
152152
{}
153153

@@ -232,20 +232,20 @@ _CCCL_DEVICE _CCCL_FORCEINLINE auto poor_apply(F&& f, Tuple&& t)
232232
-> decltype(poor_apply_impl(
233233
::cuda::std::forward<F>(f),
234234
::cuda::std::forward<Tuple>(t),
235-
::cuda::std::make_index_sequence<::cuda::std::tuple_size<::cuda::std::__libcpp_remove_reference_t<Tuple>>::value>{}))
235+
::cuda::std::make_index_sequence<::cuda::std::tuple_size<::cuda::std::remove_reference_t<Tuple>>::value>{}))
236236
{
237237
return poor_apply_impl(
238238
::cuda::std::forward<F>(f),
239239
::cuda::std::forward<Tuple>(t),
240-
::cuda::std::make_index_sequence<::cuda::std::tuple_size<::cuda::std::__libcpp_remove_reference_t<Tuple>>::value>{});
240+
::cuda::std::make_index_sequence<::cuda::std::tuple_size<::cuda::std::remove_reference_t<Tuple>>::value>{});
241241
}
242242

243243
// mult must be a power of 2
244244
template <typename Integral>
245245
_CCCL_HOST_DEVICE _CCCL_FORCEINLINE constexpr auto round_up_to_po2_multiple(Integral x, Integral mult) -> Integral
246246
{
247247
#if _CCCL_STD_VER > 2011
248-
_CCCL_ASSERT(::cuda::std::has_single_bit(static_cast<::cuda::std::__make_unsigned_t<Integral>>(mult)), "");
248+
_CCCL_ASSERT(::cuda::std::has_single_bit(static_cast<::cuda::std::make_unsigned_t<Integral>>(mult)), "");
249249
#endif // _CCCL_STD_VER > 2011
250250
return (x + mult - 1) & ~(mult - 1);
251251
}
@@ -544,15 +544,15 @@ using needs_aligned_ptr_t =
544544
>;
545545

546546
#ifdef _CUB_HAS_TRANSFORM_UBLKCP
547-
template <Algorithm Alg, typename It, ::cuda::std::__enable_if_t<needs_aligned_ptr_t<Alg>::value, int> = 0>
547+
template <Algorithm Alg, typename It, ::cuda::std::enable_if_t<needs_aligned_ptr_t<Alg>::value, int> = 0>
548548
_CCCL_DEVICE _CCCL_FORCEINLINE auto select_kernel_arg(
549549
::cuda::std::integral_constant<Algorithm, Alg>, kernel_arg<It>&& arg) -> aligned_base_ptr<value_t<It>>&&
550550
{
551551
return ::cuda::std::move(arg.aligned_ptr);
552552
}
553553
#endif // _CUB_HAS_TRANSFORM_UBLKCP
554554

555-
template <Algorithm Alg, typename It, ::cuda::std::__enable_if_t<!needs_aligned_ptr_t<Alg>::value, int> = 0>
555+
template <Algorithm Alg, typename It, ::cuda::std::enable_if_t<!needs_aligned_ptr_t<Alg>::value, int> = 0>
556556
_CCCL_DEVICE _CCCL_FORCEINLINE auto
557557
select_kernel_arg(::cuda::std::integral_constant<Algorithm, Alg>, kernel_arg<It>&& arg) -> It&&
558558
{

‎cub/cub/thread/thread_reduce.cuh

+5-5
Original file line numberDiff line numberDiff line change
@@ -131,9 +131,9 @@ ThreadReduceSequential(const Input& input, ReductionOp reduction_op)
131131
/// Specialization for DPX reduction
132132
template <typename Input, typename ReductionOp>
133133
_CCCL_NODISCARD _CCCL_DEVICE _CCCL_FORCEINLINE auto
134-
ThreadReduceDpx(const Input& input, ReductionOp reduction_op) -> ::cuda::std::__remove_cvref_t<decltype(input[0])>
134+
ThreadReduceDpx(const Input& input, ReductionOp reduction_op) -> ::cuda::std::remove_cvref_t<decltype(input[0])>
135135
{
136-
using T = ::cuda::std::__remove_cvref_t<decltype(input[0])>;
136+
using T = ::cuda::std::remove_cvref_t<decltype(input[0])>;
137137
constexpr int length = detail::static_size<Input>();
138138
T array[length];
139139
# pragma unroll
@@ -153,7 +153,7 @@ ThreadReduceDpx(const Input& input, ReductionOp reduction_op) -> ::cuda::std::__
153153
// DPX/Sequential dispatch
154154
template <typename Input,
155155
typename ReductionOp,
156-
typename ValueT = ::cuda::std::__remove_cvref_t<decltype(::cuda::std::declval<Input>()[0])>,
156+
typename ValueT = ::cuda::std::remove_cvref_t<decltype(::cuda::std::declval<Input>()[0])>,
157157
typename AccumT = ::cuda::std::__accumulator_t<ReductionOp, ValueT>,
158158
_CUB_TEMPLATE_REQUIRES(enable_dpx_reduction<Input, ReductionOp, AccumT>())>
159159
_CCCL_NODISCARD _CCCL_DEVICE _CCCL_FORCEINLINE AccumT ThreadReduce(const Input& input, ReductionOp reduction_op)
@@ -170,7 +170,7 @@ _CCCL_NODISCARD _CCCL_DEVICE _CCCL_FORCEINLINE AccumT ThreadReduce(const Input&
170170

171171
template <typename Input,
172172
typename ReductionOp,
173-
typename ValueT = ::cuda::std::__remove_cvref_t<decltype(::cuda::std::declval<Input>()[0])>,
173+
typename ValueT = ::cuda::std::remove_cvref_t<decltype(::cuda::std::declval<Input>()[0])>,
174174
typename AccumT = ::cuda::std::__accumulator_t<ReductionOp, ValueT>,
175175
_CUB_TEMPLATE_REQUIRES(!enable_dpx_reduction<Input, ReductionOp, AccumT>())>
176176
_CCCL_NODISCARD _CCCL_DEVICE _CCCL_FORCEINLINE AccumT ThreadReduce(const Input& input, ReductionOp reduction_op)
@@ -213,7 +213,7 @@ template <typename Input,
213213
typename ReductionOp,
214214
typename PrefixT,
215215
#ifndef DOXYGEN_SHOULD_SKIP_THIS // Do not document
216-
typename ValueT = ::cuda::std::__remove_cvref_t<decltype(::cuda::std::declval<Input>()[0])>,
216+
typename ValueT = ::cuda::std::remove_cvref_t<decltype(::cuda::std::declval<Input>()[0])>,
217217
#endif // !DOXYGEN_SHOULD_SKIP_THIS
218218
typename AccumT = ::cuda::std::__accumulator_t<ReductionOp, ValueT, PrefixT>>
219219
_CCCL_NODISCARD _CCCL_DEVICE _CCCL_FORCEINLINE AccumT

‎cub/cub/util_type.cuh

+1-1
Original file line numberDiff line numberDiff line change
@@ -100,7 +100,7 @@ using value_t =
100100
typename std::iterator_traits<Iterator>::value_type;
101101
# endif // defined(_CCCL_COMPILER_NVRTC)
102102

103-
template <typename It, typename FallbackT, bool = ::cuda::std::is_void<::cuda::std::__remove_pointer_t<It>>::value>
103+
template <typename It, typename FallbackT, bool = ::cuda::std::is_void<::cuda::std::remove_pointer_t<It>>::value>
104104
struct non_void_value_impl
105105
{
106106
using type = FallbackT;

‎cub/test/catch2_test_device_adjacent_difference_substract_left.cu

+1-1
Original file line numberDiff line numberDiff line change
@@ -163,7 +163,7 @@ C2H_TEST("DeviceAdjacentDifference::SubtractLeftCopy works with pointers", "[dev
163163
template <class T>
164164
struct cust_diff
165165
{
166-
template <class T2, cuda::std::__enable_if_t<cuda::std::is_same<T, T2>::value, int> = 0>
166+
template <class T2, cuda::std::enable_if_t<cuda::std::is_same<T, T2>::value, int> = 0>
167167
__host__ __device__ constexpr T2 operator()(const T2& lhs, const T2& rhs) const noexcept
168168
{
169169
return lhs - rhs;

‎cub/test/catch2_test_device_adjacent_difference_substract_right.cu

+1-1
Original file line numberDiff line numberDiff line change
@@ -92,7 +92,7 @@ C2H_TEST("DeviceAdjacentDifference::SubtractRightCopy does not change the input"
9292
template <class T>
9393
struct ref_diff
9494
{
95-
template <class T2, cuda::std::__enable_if_t<cuda::std::is_same<T, T2>::value, int> = 0>
95+
template <class T2, cuda::std::enable_if_t<cuda::std::is_same<T, T2>::value, int> = 0>
9696
__host__ __device__ constexpr T2 operator()(const T2& lhs, const T2& rhs) const noexcept
9797
{
9898
return rhs - lhs;

‎cudax/include/cuda/experimental/__async/cpos.cuh

+3-3
Original file line numberDiff line numberDiff line change
@@ -44,13 +44,13 @@ struct scheduler_t
4444
{};
4545

4646
template <class _Ty>
47-
using __sender_concept_t = typename __remove_ref_t<_Ty>::sender_concept;
47+
using __sender_concept_t = typename _CUDA_VSTD::remove_reference_t<_Ty>::sender_concept;
4848

4949
template <class _Ty>
50-
using __receiver_concept_t = typename __remove_ref_t<_Ty>::receiver_concept;
50+
using __receiver_concept_t = typename _CUDA_VSTD::remove_reference_t<_Ty>::receiver_concept;
5151

5252
template <class _Ty>
53-
using __scheduler_concept_t = typename __remove_ref_t<_Ty>::scheduler_concept;
53+
using __scheduler_concept_t = typename _CUDA_VSTD::remove_reference_t<_Ty>::scheduler_concept;
5454

5555
template <class _Ty>
5656
inline constexpr bool __is_sender = __type_valid_v<__sender_concept_t, _Ty>;

‎cudax/include/cuda/experimental/__async/meta.cuh

+12-12
Original file line numberDiff line numberDiff line change
@@ -94,7 +94,7 @@ template <class... _What>
9494
struct _ERROR : __merror_base
9595
{
9696
template <class...>
97-
using __call _LIBCUDACXX_NODEBUG_TYPE = _ERROR;
97+
using __call _CCCL_NODEBUG_ALIAS = _ERROR;
9898

9999
_ERROR operator+();
100100

@@ -138,14 +138,14 @@ template <bool _Error>
138138
struct __type_self_or_error_with_
139139
{
140140
template <class _Ty, class... _With>
141-
using __call _LIBCUDACXX_NODEBUG_TYPE = _Ty;
141+
using __call _CCCL_NODEBUG_ALIAS = _Ty;
142142
};
143143

144144
template <>
145145
struct __type_self_or_error_with_<true>
146146
{
147147
template <class _Ty, class... _With>
148-
using __call _LIBCUDACXX_NODEBUG_TYPE = decltype(__declval<_Ty&>().with(__declval<_ERROR<_With...>&>()));
148+
using __call _CCCL_NODEBUG_ALIAS = decltype(__declval<_Ty&>().with(__declval<_ERROR<_With...>&>()));
149149
};
150150

151151
template <class _Ty, class... _With>
@@ -162,7 +162,7 @@ struct __type_try__<false>
162162
using __call_q = _Fn<_Ts...>;
163163

164164
template <class _Fn, class... _Ts>
165-
using __call _LIBCUDACXX_NODEBUG_TYPE = typename _Fn::template __call<_Ts...>;
165+
using __call _CCCL_NODEBUG_ALIAS = typename _Fn::template __call<_Ts...>;
166166
};
167167

168168
template <>
@@ -172,7 +172,7 @@ struct __type_try__<true>
172172
using __call_q = __type_find_error<_Ts...>;
173173

174174
template <class _Fn, class... _Ts>
175-
using __call _LIBCUDACXX_NODEBUG_TYPE = __type_find_error<_Fn, _Ts...>;
175+
using __call _CCCL_NODEBUG_ALIAS = __type_find_error<_Fn, _Ts...>;
176176
};
177177

178178
template <class _Fn, class... _Ts>
@@ -187,7 +187,7 @@ template <class _Fn>
187187
struct __type_try
188188
{
189189
template <class... _Ts>
190-
using __call _LIBCUDACXX_NODEBUG_TYPE = __type_try_call<_Fn, _Ts...>;
190+
using __call _CCCL_NODEBUG_ALIAS = __type_try_call<_Fn, _Ts...>;
191191
};
192192

193193
template <template <class...> class _Fn, class... _Default>
@@ -198,7 +198,7 @@ template <template <class...> class _Fn>
198198
struct __type_try_quote<_Fn>
199199
{
200200
template <class... _Ts>
201-
using __call _LIBCUDACXX_NODEBUG_TYPE =
201+
using __call _CCCL_NODEBUG_ALIAS =
202202
typename __type_try__<__type_contains_error<_Ts...>>::template __call_q<_Fn, _Ts...>;
203203
};
204204

@@ -207,7 +207,7 @@ template <template <class...> class _Fn, class _Default>
207207
struct __type_try_quote<_Fn, _Default>
208208
{
209209
template <class... _Ts>
210-
using __call _LIBCUDACXX_NODEBUG_TYPE =
210+
using __call _CCCL_NODEBUG_ALIAS =
211211
typename _CUDA_VSTD::_If<__type_valid_v<_Fn, _Ts...>, //
212212
__type_try_quote<_Fn>,
213213
_CUDA_VSTD::__type_always<_Default>>::template __call<_Ts...>;
@@ -230,20 +230,20 @@ template <template <class...> class _Second, template <class...> class _First>
230230
struct __type_compose_quote
231231
{
232232
template <class... _Ts>
233-
using __call _LIBCUDACXX_NODEBUG_TYPE = _Second<_First<_Ts...>>;
233+
using __call _CCCL_NODEBUG_ALIAS = _Second<_First<_Ts...>>;
234234
};
235235

236236
struct __type_count
237237
{
238238
template <class... _Ts>
239-
using __call _LIBCUDACXX_NODEBUG_TYPE = _CUDA_VSTD::integral_constant<size_t, sizeof...(_Ts)>;
239+
using __call _CCCL_NODEBUG_ALIAS = _CUDA_VSTD::integral_constant<size_t, sizeof...(_Ts)>;
240240
};
241241

242242
template <template <class...> class _Continuation>
243243
struct __type_concat_into_quote
244244
{
245245
template <class... _Args>
246-
using __call _LIBCUDACXX_NODEBUG_TYPE =
246+
using __call _CCCL_NODEBUG_ALIAS =
247247
_CUDA_VSTD::__type_call1<_CUDA_VSTD::__type_concat<_CUDA_VSTD::__as_type_list<_Args>...>,
248248
_CUDA_VSTD::__type_quote<_Continuation>>;
249249
};
@@ -252,7 +252,7 @@ template <class _Ty>
252252
struct __type_self_or
253253
{
254254
template <class _Uy = _Ty>
255-
using __call _LIBCUDACXX_NODEBUG_TYPE = _Uy;
255+
using __call _CCCL_NODEBUG_ALIAS = _Uy;
256256
};
257257
} // namespace cuda::experimental::__async
258258

‎cudax/include/cuda/experimental/__async/type_traits.cuh

-3
Original file line numberDiff line numberDiff line change
@@ -35,9 +35,6 @@
3535
namespace cuda::experimental::__async
3636
{
3737

38-
template <class _Ty>
39-
using __remove_ref_t = _CUDA_VSTD::__libcpp_remove_reference_t<_Ty>;
40-
4138
//////////////////////////////////////////////////////////////////////////////////////////////////
4239
// __decay_t: An efficient implementation for ::std::decay
4340
#if defined(_CCCL_BUILTIN_DECAY)

‎cudax/include/cuda/experimental/__hierarchy/hierarchy_dimensions.cuh

+3-3
Original file line numberDiff line numberDiff line change
@@ -125,12 +125,12 @@ struct get_level_helper
125125

126126
template <typename QueryLevel, typename Hierarchy>
127127
_CCCL_INLINE_VAR constexpr bool has_level =
128-
detail::has_level_helper<QueryLevel, ::cuda::std::__remove_cvref_t<Hierarchy>>::value;
128+
detail::has_level_helper<QueryLevel, ::cuda::std::remove_cvref_t<Hierarchy>>::value;
129129

130130
template <typename QueryLevel, typename Hierarchy>
131131
_CCCL_INLINE_VAR constexpr bool has_level_or_unit =
132-
detail::has_level_helper<QueryLevel, ::cuda::std::__remove_cvref_t<Hierarchy>>::value
133-
|| detail::has_unit<QueryLevel, ::cuda::std::__remove_cvref_t<Hierarchy>>::value;
132+
detail::has_level_helper<QueryLevel, ::cuda::std::remove_cvref_t<Hierarchy>>::value
133+
|| detail::has_unit<QueryLevel, ::cuda::std::remove_cvref_t<Hierarchy>>::value;
134134

135135
namespace detail
136136
{

‎docs/repo.toml

+1-1
Original file line numberDiff line numberDiff line change
@@ -436,7 +436,7 @@ doxygen_predefined = [
436436
"_LIBCUDACXX_AND=&&",
437437
"_LIBCUDACXX_EAT_REST(x)=",
438438
"_LIBCUDACXX_GLOBAL_CONSTANT=inline",
439-
"_LIBCUDACXX_REQUIRES(x)= ::cuda::std::__enable_if_t<x, int> = 0>",
439+
"_LIBCUDACXX_REQUIRES(x)= ::cuda::std::enable_if_t<x, int> = 0>",
440440
"_LIBCUDACXX_TEMPLATE(x)=template<x, ",
441441
"_LIBCUDACXX_TRAILING_REQUIRES(x)=-> x _LIBCUDACXX_EAT_REST",
442442
"LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE=",

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

+9-9
Original file line numberDiff line numberDiff line change
@@ -39,12 +39,12 @@ _LIBCUDACXX_BEGIN_NAMESPACE_CUDA
3939
//! @pre \p __b must be positive
4040
template <class _Tp,
4141
class _Up,
42-
_CUDA_VSTD::__enable_if_t<_CCCL_TRAIT(_CUDA_VSTD::is_unsigned, _Tp), int> = 0,
43-
_CUDA_VSTD::__enable_if_t<_CCCL_TRAIT(_CUDA_VSTD::is_integral, _Up), int> = 0>
42+
_CUDA_VSTD::enable_if_t<_CCCL_TRAIT(_CUDA_VSTD::is_unsigned, _Tp), int> = 0,
43+
_CUDA_VSTD::enable_if_t<_CCCL_TRAIT(_CUDA_VSTD::is_integral, _Up), int> = 0>
4444
_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 _Tp ceil_div(const _Tp __a, const _Up __b) noexcept
4545
{
4646
_CCCL_ASSERT(__b > _Up(0), "cuda::ceil_div: b must be positive");
47-
using _UCommon = _CUDA_VSTD::__make_unsigned_t<_CUDA_VSTD::__common_type_t<_Tp, _Up>>;
47+
using _UCommon = _CUDA_VSTD::make_unsigned_t<_CUDA_VSTD::common_type_t<_Tp, _Up>>;
4848
const auto __res = static_cast<_UCommon>(__a) / static_cast<_UCommon>(__b);
4949
return static_cast<_Tp>(__res + (__res * static_cast<_UCommon>(__b) != static_cast<_UCommon>(__a)));
5050
}
@@ -56,13 +56,13 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 _Tp ceil_div(con
5656
//! @pre \p __b must be positive
5757
template <class _Tp,
5858
class _Up,
59-
_CUDA_VSTD::__enable_if_t<_CCCL_TRAIT(_CUDA_VSTD::is_signed, _Tp), int> = 0,
60-
_CUDA_VSTD::__enable_if_t<_CCCL_TRAIT(_CUDA_VSTD::is_integral, _Up), int> = 0>
59+
_CUDA_VSTD::enable_if_t<_CCCL_TRAIT(_CUDA_VSTD::is_signed, _Tp), int> = 0,
60+
_CUDA_VSTD::enable_if_t<_CCCL_TRAIT(_CUDA_VSTD::is_integral, _Up), int> = 0>
6161
_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 _Tp ceil_div(const _Tp __a, const _Up __b) noexcept
6262
{
6363
_CCCL_ASSERT(__a >= _Tp(0), "cuda::ceil_div: a must be non negative");
6464
_CCCL_ASSERT(__b > _Up(0), "cuda::ceil_div: b must be positive");
65-
using _UCommon = _CUDA_VSTD::__make_unsigned_t<_CUDA_VSTD::__common_type_t<_Tp, _Up>>;
65+
using _UCommon = _CUDA_VSTD::make_unsigned_t<_CUDA_VSTD::common_type_t<_Tp, _Up>>;
6666
// Due to the precondition `__a >= 0` we can safely cast to unsigned without danger of overflowing
6767
return static_cast<_Tp>((static_cast<_UCommon>(__a) + static_cast<_UCommon>(__b) - 1) / static_cast<_UCommon>(__b));
6868
}
@@ -74,11 +74,11 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 _Tp ceil_div(con
7474
//! @pre \p __b must be positive
7575
template <class _Tp,
7676
class _Up,
77-
_CUDA_VSTD::__enable_if_t<_CCCL_TRAIT(_CUDA_VSTD::is_integral, _Tp), int> = 0,
78-
_CUDA_VSTD::__enable_if_t<_CCCL_TRAIT(_CUDA_VSTD::is_enum, _Up), int> = 0>
77+
_CUDA_VSTD::enable_if_t<_CCCL_TRAIT(_CUDA_VSTD::is_integral, _Tp), int> = 0,
78+
_CUDA_VSTD::enable_if_t<_CCCL_TRAIT(_CUDA_VSTD::is_enum, _Up), int> = 0>
7979
_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 _Tp ceil_div(const _Tp __a, const _Up __b) noexcept
8080
{
81-
return ::cuda::ceil_div(__a, static_cast<_CUDA_VSTD::__underlying_type_t<_Up>>(__b));
81+
return ::cuda::ceil_div(__a, static_cast<_CUDA_VSTD::underlying_type_t<_Up>>(__b));
8282
}
8383

8484
_LIBCUDACXX_END_NAMESPACE_CUDA

‎libcudacxx/include/cuda/__functional/maximum.h

+1-1
Original file line numberDiff line numberDiff line change
@@ -42,7 +42,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT maximum<void>
4242
{
4343
_CCCL_EXEC_CHECK_DISABLE
4444
template <class _T1, class _T2>
45-
_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr _CUDA_VSTD::__common_type_t<_T1, _T2>
45+
_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr _CUDA_VSTD::common_type_t<_T1, _T2>
4646
operator()(const _T1& __lhs, const _T2& __rhs) const noexcept(noexcept((__lhs < __rhs) ? __rhs : __lhs))
4747
{
4848
return (__lhs < __rhs) ? __rhs : __lhs;

‎libcudacxx/include/cuda/__functional/minimum.h

+1-1
Original file line numberDiff line numberDiff line change
@@ -42,7 +42,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT minimum<void>
4242
{
4343
_CCCL_EXEC_CHECK_DISABLE
4444
template <class _T1, class _T2>
45-
_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr _CUDA_VSTD::__common_type_t<_T1, _T2>
45+
_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr _CUDA_VSTD::common_type_t<_T1, _T2>
4646
operator()(const _T1& __lhs, const _T2& __rhs) const noexcept(noexcept((__lhs < __rhs) ? __lhs : __rhs))
4747
{
4848
return (__lhs < __rhs) ? __lhs : __rhs;

0 commit comments

Comments
 (0)