From 6cbaa6ee03d87f68515896e002e34a884c7d13c4 Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Fri, 26 Apr 2024 18:17:01 +0200 Subject: [PATCH] Move `std::midpoint` to its own file --- .../include/cuda/std/__numeric/midpoint.h | 104 ++++++++++++ .../cuda/std/__type_traits/is_null_pointer.h | 2 - .../cuda/std/detail/libcxx/include/numeric | 81 +--------- .../midpoint.float.pass.cpp | 130 +++++++++++++++ .../midpoint.integer.pass.cpp | 151 ++++++++++++++++++ .../midpoint.pointer.pass.cpp | 91 +++++++++++ .../numeric.ops.midpoint/midpoint.verify.cpp | 44 +++++ libcudacxx/test/support/fp_compare.h | 17 +- 8 files changed, 531 insertions(+), 89 deletions(-) create mode 100644 libcudacxx/include/cuda/std/__numeric/midpoint.h create mode 100644 libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.midpoint/midpoint.float.pass.cpp create mode 100644 libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.midpoint/midpoint.integer.pass.cpp create mode 100644 libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.midpoint/midpoint.pointer.pass.cpp create mode 100644 libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.midpoint/midpoint.verify.cpp diff --git a/libcudacxx/include/cuda/std/__numeric/midpoint.h b/libcudacxx/include/cuda/std/__numeric/midpoint.h new file mode 100644 index 00000000000..385c851fde4 --- /dev/null +++ b/libcudacxx/include/cuda/std/__numeric/midpoint.h @@ -0,0 +1,104 @@ +// -*- C++ -*- +//===----------------------------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +#ifndef _LIBCUDACXX___NUMERIC_MIDPOINT_H +#define _LIBCUDACXX___NUMERIC_MIDPOINT_H + +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +// comes last +#include + +_LIBCUDACXX_BEGIN_NAMESPACE_STD + +template +_CCCL_NODISCARD _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 + __enable_if_t<_LIBCUDACXX_TRAIT(is_integral, _Tp) && !_LIBCUDACXX_TRAIT(is_same, bool, _Tp) + && !_LIBCUDACXX_TRAIT(is_null_pointer, _Tp), + _Tp> + midpoint(_Tp __a, _Tp __b) noexcept +{ + using _Up = __make_unsigned_t<_Tp>; + + if (__a > __b) + { + const _Up __diff = _Up(__a) - _Up(__b); + return static_cast<_Tp>(__a - static_cast<_Tp>(__diff / 2)); + } + else + { + const _Up __diff = _Up(__b) - _Up(__a); + return static_cast<_Tp>(__a + static_cast<_Tp>(__diff / 2)); + } +} + +template < + class _Tp, + __enable_if_t<_LIBCUDACXX_TRAIT(is_object, _Tp) && !_LIBCUDACXX_TRAIT(is_void, _Tp) && (sizeof(_Tp) > 0), int> = 0> +_CCCL_NODISCARD _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 _Tp* midpoint(_Tp* __a, _Tp* __b) noexcept +{ + return __a + _CUDA_VSTD::midpoint(ptrdiff_t(0), __b - __a); +} + +template +_CCCL_NODISCARD _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 int __sign(_Tp __val) +{ + return (_Tp(0) < __val) - (__val < _Tp(0)); +} + +template +_CCCL_NODISCARD _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 _Fp __fp_abs(_Fp __f) +{ + return __f >= 0 ? __f : -__f; +} + +template +_CCCL_NODISCARD _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 + __enable_if_t<_LIBCUDACXX_TRAIT(is_floating_point, _Fp), _Fp> + midpoint(_Fp __a, _Fp __b) noexcept +{ + _CCCL_CONSTEXPR_CXX14 _Fp __lo = numeric_limits<_Fp>::min() * 2; + _CCCL_CONSTEXPR_CXX14 _Fp __hi = numeric_limits<_Fp>::max() / 2; + return _CUDA_VSTD::__fp_abs(__a) <= __hi && _CUDA_VSTD::__fp_abs(__b) <= __hi + ? // typical case: overflow is impossible + (__a + __b) / 2 + : // always correctly rounded + _CUDA_VSTD::__fp_abs(__a) < __lo ? __a + __b / 2 : // not safe to halve a + _CUDA_VSTD::__fp_abs(__b) < __lo + ? __a / 2 + __b + : // not safe to halve b + __a / 2 + __b / 2; // otherwise correctly rounded +} + +_LIBCUDACXX_END_NAMESPACE_STD + +#include + +#endif // _LIBCUDACXX___NUMERIC_MIDPOINT_H diff --git a/libcudacxx/include/cuda/std/__type_traits/is_null_pointer.h b/libcudacxx/include/cuda/std/__type_traits/is_null_pointer.h index 9b9ed837d81..46f2f5ca847 100644 --- a/libcudacxx/include/cuda/std/__type_traits/is_null_pointer.h +++ b/libcudacxx/include/cuda/std/__type_traits/is_null_pointer.h @@ -37,11 +37,9 @@ template struct _LIBCUDACXX_TEMPLATE_VIS __is_nullptr_t : public __is_nullptr_t_impl<__remove_cv_t<_Tp>> {}; -#if _CCCL_STD_VER > 2011 template struct _LIBCUDACXX_TEMPLATE_VIS is_null_pointer : public __is_nullptr_t_impl<__remove_cv_t<_Tp>> {}; -#endif #if _CCCL_STD_VER > 2011 && !defined(_LIBCUDACXX_HAS_NO_VARIABLE_TEMPLATES) template diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/numeric b/libcudacxx/include/cuda/std/detail/libcxx/include/numeric index e128bf0a05f..ebb27337e26 100644 --- a/libcudacxx/include/cuda/std/detail/libcxx/include/numeric +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/numeric @@ -155,7 +155,6 @@ floating_point midpoint(floating_point a, floating_point b); // C++20 # pragma system_header #endif // no system header -#include #include #include #include @@ -163,89 +162,15 @@ floating_point midpoint(floating_point a, floating_point b); // C++20 #include #include #include +#include #include #include #include #include #include -#include -#include // for isnormal #include // all public C++ headers provide the assertion handler -#include -#include -#include // for numeric_limits -#include - -_LIBCUDACXX_BEGIN_NAMESPACE_STD - -#ifndef __cuda_std__ -# if _CCCL_STD_VER > 2017 -template -_LIBCUDACXX_INLINE_VISIBILITY constexpr enable_if_t< - is_integral_v<_Tp> && !is_same_v && !is_null_pointer_v<_Tp>, - _Tp> -midpoint(_Tp __a, _Tp __b) noexcept _LIBCUDACXX_DISABLE_UBSAN_UNSIGNED_INTEGER_CHECK -{ - using _Up = std::make_unsigned_t<_Tp>; - - int __sign = 1; - _Up __m = __a; - _Up __M = __b; - if (__a > __b) - { - __sign = -1; - __m = __b; - __M = __a; - } - return __a + __sign * _Tp(_Up(__M - __m) >> 1); -} - -template -_LIBCUDACXX_INLINE_VISIBILITY constexpr enable_if_t< - is_pointer_v<_TPtr> && is_object_v> && !is_void_v> - && (sizeof(remove_pointer_t<_TPtr>) > 0), - _TPtr> -midpoint(_TPtr __a, _TPtr __b) noexcept -{ - return __a + _CUDA_VSTD::midpoint(ptrdiff_t(0), __b - __a); -} - -template -constexpr int __sign(_Tp __val) -{ - return (_Tp(0) < __val) - (__val < _Tp(0)); -} -template -constexpr _Fp __fp_abs(_Fp __f) -{ - return __f >= 0 ? __f : -__f; -} - -template -_LIBCUDACXX_INLINE_VISIBILITY constexpr enable_if_t, _Fp> midpoint(_Fp __a, _Fp __b) noexcept -{ - constexpr _Fp __lo = numeric_limits<_Fp>::min() * 2; - constexpr _Fp __hi = numeric_limits<_Fp>::max() / 2; - return __fp_abs(__a) <= __hi && __fp_abs(__b) <= __hi - ? // typical case: overflow is impossible - (__a + __b) / 2 - : // always correctly rounded - __fp_abs(__a) < __lo ? __a + __b / 2 : // not safe to halve a - __fp_abs(__a) < __lo ? __a / 2 + __b - : // not safe to halve b - __a / 2 + __b / 2; // otherwise correctly rounded -} - -# endif // _CCCL_STD_VER > 2017 -#endif // __cuda_std__ - -_LIBCUDACXX_END_NAMESPACE_STD - -#if defined(_LIBCUDACXX_HAS_PARALLEL_ALGORITHMS) && _CCCL_STD_VER >= 2017 -# include <__pstl_numeric> -#endif - -#include //__cuda_std__ +// standard mandated include +#include #endif // _LIBCUDACXX_NUMERIC diff --git a/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.midpoint/midpoint.float.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.midpoint/midpoint.float.pass.cpp new file mode 100644 index 00000000000..a2374452993 --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.midpoint/midpoint.float.pass.cpp @@ -0,0 +1,130 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// + +// template +// _Tp midpoint(_Float __a, _Float __b) noexcept +// + +#include +#include +#include + +#include "fp_compare.h" +#include "test_macros.h" + +// Totally arbitrary picks for precision +template +__host__ __device__ constexpr T fp_error_pct(); + +template <> +__host__ __device__ constexpr float fp_error_pct() +{ + return 1.0e-4f; +} + +template <> +__host__ __device__ constexpr double fp_error_pct() +{ + return 1.0e-12; +} + +template <> +__host__ __device__ constexpr long double fp_error_pct() +{ + return 1.0e-13l; +} + +template +__host__ __device__ void fp_test() +{ + ASSERT_SAME_TYPE(T, decltype(cuda::std::midpoint(T(), T()))); + ASSERT_NOEXCEPT(cuda::std::midpoint(T(), T())); + + constexpr T maxV = cuda::std::numeric_limits::max(); + constexpr T minV = cuda::std::numeric_limits::min(); + + // Things that can be compared exactly + assert((cuda::std::midpoint(T(0), T(0)) == T(0))); + assert((cuda::std::midpoint(T(2), T(4)) == T(3))); + assert((cuda::std::midpoint(T(4), T(2)) == T(3))); + assert((cuda::std::midpoint(T(3), T(4)) == T(3.5))); + assert((cuda::std::midpoint(T(0), T(0.4)) == T(0.2))); + + // Things that can't be compared exactly + constexpr T pct = fp_error_pct(); + assert((fptest_close_pct(cuda::std::midpoint(T(1.3), T(11.4)), T(6.35), pct))); + assert((fptest_close_pct(cuda::std::midpoint(T(11.33), T(31.45)), T(21.39), pct))); + assert((fptest_close_pct(cuda::std::midpoint(T(-1.3), T(11.4)), T(5.05), pct))); + assert((fptest_close_pct(cuda::std::midpoint(T(11.4), T(-1.3)), T(5.05), pct))); + assert((fptest_close_pct(cuda::std::midpoint(T(0.1), T(0.4)), T(0.25), pct))); + + assert((fptest_close_pct(cuda::std::midpoint(T(11.2345), T(14.5432)), T(12.88885), pct))); + + // From e to pi + assert((fptest_close_pct(cuda::std::midpoint(T(2.71828182845904523536028747135266249775724709369995), + T(3.14159265358979323846264338327950288419716939937510)), + T(2.92993724102441923691146542731608269097720824653752), + pct))); + + assert((fptest_close_pct(cuda::std::midpoint(maxV, T(0)), maxV / 2, pct))); + assert((fptest_close_pct(cuda::std::midpoint(T(0), maxV), maxV / 2, pct))); + assert((fptest_close_pct(cuda::std::midpoint(minV, T(0)), minV / 2, pct))); + assert((fptest_close_pct(cuda::std::midpoint(T(0), minV), minV / 2, pct))); + assert((fptest_close_pct(cuda::std::midpoint(maxV, maxV), maxV, pct))); + assert((fptest_close_pct(cuda::std::midpoint(minV, minV), minV, pct))); + assert((fptest_close_pct(cuda::std::midpoint(maxV, minV), maxV / 2, pct))); + assert((fptest_close_pct(cuda::std::midpoint(minV, maxV), maxV / 2, pct))); + + // Near the min and the max + assert((fptest_close_pct(cuda::std::midpoint(maxV * T(0.75), maxV * T(0.50)), maxV * T(0.625), pct))); + assert((fptest_close_pct(cuda::std::midpoint(maxV * T(0.50), maxV * T(0.75)), maxV * T(0.625), pct))); + assert((fptest_close_pct(cuda::std::midpoint(minV * T(2), minV * T(8)), minV * T(5), pct))); + + // Big numbers of different signs + assert((fptest_close_pct(cuda::std::midpoint(maxV * T(0.75), maxV * T(-0.5)), maxV * T(0.125), pct))); + assert((fptest_close_pct(cuda::std::midpoint(maxV * T(-0.75), maxV * T(0.5)), maxV * T(-0.125), pct))); + +#if !defined(TEST_COMPILER_NVRTC) // missing nextafter + // Check two values "close to each other" + T d1 = T(3.14); + T d0 = cuda::std::nextafter(d1, T(2)); + T d2 = cuda::std::nextafter(d1, T(5)); + assert(d0 < d1); // sanity checking + assert(d1 < d2); // sanity checking + + // Since there's nothing in between, the midpoint has to be one or the other + T res; + res = cuda::std::midpoint(d0, d1); + assert(res == d0 || res == d1); + assert(d0 <= res); + assert(res <= d1); + res = cuda::std::midpoint(d1, d0); + assert(res == d0 || res == d1); + assert(d0 <= res); + assert(res <= d1); + + res = cuda::std::midpoint(d1, d2); + assert(res == d1 || res == d2); + assert(d1 <= res); + assert(res <= d2); + res = cuda::std::midpoint(d2, d1); + assert(res == d1 || res == d2); + assert(d1 <= res); + assert(res <= d2); +#endif // !TEST_COMPILER_NVRTC +} + +int main(int, char**) +{ + fp_test(); + fp_test(); + + return 0; +} diff --git a/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.midpoint/midpoint.integer.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.midpoint/midpoint.integer.pass.cpp new file mode 100644 index 00000000000..fdd81e9a69c --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.midpoint/midpoint.integer.pass.cpp @@ -0,0 +1,151 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// +// +// + +// template +// _Tp midpoint(_Tp __a, _Tp __b) noexcept +// + +#include +#include +#include +#include +#include + +#include "test_macros.h" + +template +__host__ __device__ TEST_CONSTEXPR_CXX14 void signed_test() +{ + constexpr T zero{0}; + constexpr T one{1}; + constexpr T two{2}; + constexpr T three{3}; + constexpr T four{4}; + + ASSERT_SAME_TYPE(decltype(cuda::std::midpoint(T(), T())), T); + ASSERT_NOEXCEPT(cuda::std::midpoint(T(), T())); + using limits = cuda::std::numeric_limits; + + assert(cuda::std::midpoint(one, three) == two); + assert(cuda::std::midpoint(three, one) == two); + + assert(cuda::std::midpoint(zero, zero) == zero); + assert(cuda::std::midpoint(zero, two) == one); + assert(cuda::std::midpoint(two, zero) == one); + assert(cuda::std::midpoint(two, two) == two); + + assert(cuda::std::midpoint(one, four) == two); + assert(cuda::std::midpoint(four, one) == three); + assert(cuda::std::midpoint(three, four) == three); + assert(cuda::std::midpoint(four, three) == four); + + assert(cuda::std::midpoint(T(3), T(4)) == T(3)); + assert(cuda::std::midpoint(T(4), T(3)) == T(4)); + assert(cuda::std::midpoint(T(-3), T(4)) == T(0)); + assert(cuda::std::midpoint(T(-4), T(3)) == T(-1)); + assert(cuda::std::midpoint(T(3), T(-4)) == T(0)); + assert(cuda::std::midpoint(T(4), T(-3)) == T(1)); + assert(cuda::std::midpoint(T(-3), T(-4)) == T(-3)); + assert(cuda::std::midpoint(T(-4), T(-3)) == T(-4)); + + assert(cuda::std::midpoint(limits::min(), limits::max()) == T(-1)); + assert(cuda::std::midpoint(limits::max(), limits::min()) == T(0)); + + assert(cuda::std::midpoint(limits::min(), T(6)) == limits::min() / 2 + 3); + assert(cuda::std::midpoint(T(6), limits::min()) == limits::min() / 2 + 3); + assert(cuda::std::midpoint(limits::max(), T(6)) == limits::max() / 2 + 4); + assert(cuda::std::midpoint(T(6), limits::max()) == limits::max() / 2 + 3); + + assert(cuda::std::midpoint(limits::min(), T(-6)) == limits::min() / 2 - 3); + assert(cuda::std::midpoint(T(-6), limits::min()) == limits::min() / 2 - 3); + assert(cuda::std::midpoint(limits::max(), T(-6)) == limits::max() / 2 - 2); + assert(cuda::std::midpoint(T(-6), limits::max()) == limits::max() / 2 - 3); +} + +template +__host__ __device__ TEST_CONSTEXPR_CXX14 void unsigned_test() +{ + constexpr T zero{0}; + constexpr T one{1}; + constexpr T two{2}; + constexpr T three{3}; + constexpr T four{4}; + + ASSERT_SAME_TYPE(decltype(cuda::std::midpoint(T(), T())), T); + ASSERT_NOEXCEPT(cuda::std::midpoint(T(), T())); + using limits = cuda::std::numeric_limits; + const T half_way = (limits::max() - limits::min()) / 2; + + assert(cuda::std::midpoint(one, three) == two); + assert(cuda::std::midpoint(three, one) == two); + + assert(cuda::std::midpoint(zero, zero) == zero); + assert(cuda::std::midpoint(zero, two) == one); + assert(cuda::std::midpoint(two, zero) == one); + assert(cuda::std::midpoint(two, two) == two); + + assert(cuda::std::midpoint(one, four) == two); + assert(cuda::std::midpoint(four, one) == three); + assert(cuda::std::midpoint(three, four) == three); + assert(cuda::std::midpoint(four, three) == four); + + assert(cuda::std::midpoint(limits::min(), limits::max()) == T(half_way)); + assert(cuda::std::midpoint(limits::max(), limits::min()) == T(half_way + 1)); + + assert(cuda::std::midpoint(limits::min(), T(6)) == limits::min() / 2 + 3); + assert(cuda::std::midpoint(T(6), limits::min()) == limits::min() / 2 + 3); + assert(cuda::std::midpoint(limits::max(), T(6)) == half_way + 4); + assert(cuda::std::midpoint(T(6), limits::max()) == half_way + 3); +} + +__host__ __device__ TEST_CONSTEXPR_CXX14 bool test() +{ + signed_test(); + signed_test(); + signed_test(); + signed_test(); + signed_test(); + + signed_test(); + signed_test(); + signed_test(); + signed_test(); + + unsigned_test(); + unsigned_test(); + unsigned_test(); + unsigned_test(); + unsigned_test(); + + unsigned_test(); + unsigned_test(); + unsigned_test(); + unsigned_test(); + +#ifndef TEST_HAS_NO_INT128_T + unsigned_test<__uint128_t>(); + signed_test<__int128_t>(); +#endif // !TEST_HAS_NO_INT128_T + + // int_test(); + signed_test(); + unsigned_test(); + + return true; +} + +int main(int, char**) +{ + test(); +#if TEST_STD_VER >= 2014 + static_assert(test(), ""); +#endif // TEST_STD_VER >= 2014 + return 0; +} diff --git a/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.midpoint/midpoint.pointer.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.midpoint/midpoint.pointer.pass.cpp new file mode 100644 index 00000000000..aca7958684e --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.midpoint/midpoint.pointer.pass.cpp @@ -0,0 +1,91 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// +// +// + +// template +// _Tp* midpoint(_Tp* __a, _Tp* __b) noexcept +// + +#include +#include + +#include "test_macros.h" + +#if TEST_STD_VER >= 2014 +template +__host__ __device__ TEST_CONSTEXPR_CXX14 void constexpr_test() +{ + constexpr T array[1000] = {}; + ASSERT_SAME_TYPE(decltype(cuda::std::midpoint(array, array)), const T*); + ASSERT_NOEXCEPT(cuda::std::midpoint(array, array)); + + static_assert(cuda::std::midpoint(array, array) == array, ""); + static_assert(cuda::std::midpoint(array, array + 1000) == array + 500, ""); + + static_assert(cuda::std::midpoint(array, array + 9) == array + 4, ""); + static_assert(cuda::std::midpoint(array, array + 10) == array + 5, ""); + static_assert(cuda::std::midpoint(array, array + 11) == array + 5, ""); + static_assert(cuda::std::midpoint(array + 9, array) == array + 5, ""); + static_assert(cuda::std::midpoint(array + 10, array) == array + 5, ""); + static_assert(cuda::std::midpoint(array + 11, array) == array + 6, ""); +} +#endif // TEST_STD_VER >= 2014 + +template +__host__ __device__ void runtime_test() +{ + T array[1000] = {}; // we need an array to make valid pointers + ASSERT_SAME_TYPE(decltype(cuda::std::midpoint(array, array)), T*); + ASSERT_NOEXCEPT(cuda::std::midpoint(array, array)); + + assert(cuda::std::midpoint(array, array) == array); + assert(cuda::std::midpoint(array, array + 1000) == array + 500); + + assert(cuda::std::midpoint(array, array + 9) == array + 4); + assert(cuda::std::midpoint(array, array + 10) == array + 5); + assert(cuda::std::midpoint(array, array + 11) == array + 5); + assert(cuda::std::midpoint(array + 9, array) == array + 5); + assert(cuda::std::midpoint(array + 10, array) == array + 5); + assert(cuda::std::midpoint(array + 11, array) == array + 6); + + // explicit instantiation + ASSERT_SAME_TYPE(decltype(cuda::std::midpoint(array, array)), T*); + ASSERT_NOEXCEPT(cuda::std::midpoint(array, array)); + assert(cuda::std::midpoint(array, array) == array); + assert(cuda::std::midpoint(array, array + 1000) == array + 500); +} + +template +__host__ __device__ void pointer_test() +{ + runtime_test(); + runtime_test(); + runtime_test(); + runtime_test(); + + // The constexpr tests are always const, but we can test them anyway. +#if TEST_STD_VER >= 2014 && !defined(TEST_COMPILER_CUDACC_BELOW_11_3) + constexpr_test(); + constexpr_test(); + +# if !defined(TEST_COMPILER_GCC) + constexpr_test(); + constexpr_test(); +# endif // !TEST_COMPILER_GCC +#endif // TEST_STD_VER >= 2014 && !TEST_COMPILER_CUDACC_BELOW_11_3 +} + +int main(int, char**) +{ + pointer_test(); + pointer_test(); + pointer_test(); + + return 0; +} diff --git a/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.midpoint/midpoint.verify.cpp b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.midpoint/midpoint.verify.cpp new file mode 100644 index 00000000000..aa67b8fa1e8 --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.midpoint/midpoint.verify.cpp @@ -0,0 +1,44 @@ +//===----------------------------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// +// +// + +// template +// _Tp midpoint(_Tp __a, _Tp __b) noexcept + +// An overload exists for each of char and all arithmetic types except bool. + +#include + +#include "test_macros.h" + +__host__ __device__ int func1() +{ + return 1; +} +__host__ __device__ int func2() +{ + return 2; +} + +struct Incomplete; +Incomplete* ip = nullptr; +void* vp = nullptr; + +int main(int, char**) +{ + (void) cuda::std::midpoint(false, true); // expected-error {{no matching function for call to 'midpoint'}} + + // A couple of odd pointer types that should fail + (void) cuda::std::midpoint(nullptr, nullptr); // expected-error {{no matching function for call to 'midpoint'}} + (void) cuda::std::midpoint(func1, func2); // expected-error {{no matching function for call to 'midpoint'}} + (void) cuda::std::midpoint(ip, ip); // expected-error {{no matching function for call to 'midpoint'}} + (void) cuda::std::midpoint(vp, vp); // expected-error {{no matching function for call to 'midpoint'}} + + return 0; +} diff --git a/libcudacxx/test/support/fp_compare.h b/libcudacxx/test/support/fp_compare.h index 2e6b4c0d661..fee0efa7cd2 100644 --- a/libcudacxx/test/support/fp_compare.h +++ b/libcudacxx/test/support/fp_compare.h @@ -8,16 +8,15 @@ #ifndef SUPPORT_FP_COMPARE_H #define SUPPORT_FP_COMPARE_H +#include // for cuda::std::max #include - -#include // for std::max -#include // for std::abs +#include // for cuda::std::abs // See // https://www.boost.org/doc/libs/1_70_0/libs/test/doc/html/boost_test/testing_tools/extended_comparison/floating_point/floating_points_comparison_theory.html template -bool fptest_close(T val, T expected, T eps) +__host__ __device__ bool fptest_close(T val, T expected, T eps) { constexpr T zero = T(0); assert(eps >= zero); @@ -29,18 +28,18 @@ bool fptest_close(T val, T expected, T eps) } if (val == zero) { - return std::abs(expected) <= eps; + return cuda::std::abs(expected) <= eps; } if (expected == zero) { - return std::abs(val) <= eps; + return cuda::std::abs(val) <= eps; } - return std::abs(val - expected) < eps && std::abs(val - expected) / std::abs(val) < eps; + return cuda::std::abs(val - expected) < eps && cuda::std::abs(val - expected) / cuda::std::abs(val) < eps; } template -bool fptest_close_pct(T val, T expected, T percent) +__host__ __device__ bool fptest_close_pct(T val, T expected, T percent) { constexpr T zero = T(0); assert(percent >= zero); @@ -50,7 +49,7 @@ bool fptest_close_pct(T val, T expected, T percent) { return val == expected; } - T eps = (percent / T(100)) * std::max(std::abs(val), std::abs(expected)); + T eps = (percent / T(100)) * cuda::std::max(cuda::std::abs(val), cuda::std::abs(expected)); return fptest_close(val, expected, eps); }