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 <cuda/std/detail/__config>
+
+#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 <cuda/std/__type_traits/enable_if.h>
+#include <cuda/std/__type_traits/is_floating_point.h>
+#include <cuda/std/__type_traits/is_integral.h>
+#include <cuda/std/__type_traits/is_null_pointer.h>
+#include <cuda/std/__type_traits/is_object.h>
+#include <cuda/std/__type_traits/is_pointer.h>
+#include <cuda/std/__type_traits/is_same.h>
+#include <cuda/std/__type_traits/is_void.h>
+#include <cuda/std/__type_traits/make_unsigned.h>
+#include <cuda/std/__type_traits/remove_pointer.h>
+#include <cuda/std/cstddef>
+#include <cuda/std/limits>
+
+// comes last
+#include <cuda/std/detail/__pragma_push>
+
+_LIBCUDACXX_BEGIN_NAMESPACE_STD
+
+template <class _Tp>
+_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 <typename _Tp>
+_CCCL_NODISCARD _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 int __sign(_Tp __val)
+{
+  return (_Tp(0) < __val) - (__val < _Tp(0));
+}
+
+template <typename _Fp>
+_CCCL_NODISCARD _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 _Fp __fp_abs(_Fp __f)
+{
+  return __f >= 0 ? __f : -__f;
+}
+
+template <class _Fp>
+_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 <cuda/std/detail/__pragma_pop>
+
+#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 <class _Tp>
 struct _LIBCUDACXX_TEMPLATE_VIS __is_nullptr_t : public __is_nullptr_t_impl<__remove_cv_t<_Tp>>
 {};
 
-#if _CCCL_STD_VER > 2011
 template <class _Tp>
 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 <class _Tp>
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 <cuda/std/__iterator/iterator_traits.h>
 #include <cuda/std/__numeric/accumulate.h>
 #include <cuda/std/__numeric/adjacent_difference.h>
 #include <cuda/std/__numeric/exclusive_scan.h>
@@ -163,89 +162,15 @@ floating_point  midpoint(floating_point a, floating_point b);    // C++20
 #include <cuda/std/__numeric/inclusive_scan.h>
 #include <cuda/std/__numeric/inner_product.h>
 #include <cuda/std/__numeric/iota.h>
+#include <cuda/std/__numeric/midpoint.h>
 #include <cuda/std/__numeric/partial_sum.h>
 #include <cuda/std/__numeric/reduce.h>
 #include <cuda/std/__numeric/transform_exclusive_scan.h>
 #include <cuda/std/__numeric/transform_inclusive_scan.h>
 #include <cuda/std/__numeric/transform_reduce.h>
-#include <cuda/std/__utility/move.h>
-#include <cuda/std/cmath> // for isnormal
 #include <cuda/std/detail/libcxx/include/__assert> // all public C++ headers provide the assertion handler
-#include <cuda/std/detail/libcxx/include/__pragma_push>
-#include <cuda/std/functional>
-#include <cuda/std/limits> // for numeric_limits
-#include <cuda/std/version>
-
-_LIBCUDACXX_BEGIN_NAMESPACE_STD
-
-#ifndef __cuda_std__
-#  if _CCCL_STD_VER > 2017
-template <class _Tp>
-_LIBCUDACXX_INLINE_VISIBILITY constexpr enable_if_t<
-  is_integral_v<_Tp> && !is_same_v<bool, _Tp> && !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 <class _TPtr>
-_LIBCUDACXX_INLINE_VISIBILITY constexpr enable_if_t<
-  is_pointer_v<_TPtr> && is_object_v<remove_pointer_t<_TPtr>> && !is_void_v<remove_pointer_t<_TPtr>>
-    && (sizeof(remove_pointer_t<_TPtr>) > 0),
-  _TPtr>
-midpoint(_TPtr __a, _TPtr __b) noexcept
-{
-  return __a + _CUDA_VSTD::midpoint(ptrdiff_t(0), __b - __a);
-}
-
-template <typename _Tp>
-constexpr int __sign(_Tp __val)
-{
-  return (_Tp(0) < __val) - (__val < _Tp(0));
-}
 
-template <typename _Fp>
-constexpr _Fp __fp_abs(_Fp __f)
-{
-  return __f >= 0 ? __f : -__f;
-}
-
-template <class _Fp>
-_LIBCUDACXX_INLINE_VISIBILITY constexpr enable_if_t<is_floating_point_v<_Fp>, _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/detail/libcxx/include/__pragma_pop> //__cuda_std__
+// standard mandated include
+#include <cuda/std/version>
 
 #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
+//
+//===----------------------------------------------------------------------===//
+
+// <numeric>
+
+// template <class _Float>
+// _Tp midpoint(_Float __a, _Float __b) noexcept
+//
+
+#include <cuda/std/__numeric_>
+#include <cuda/std/cassert>
+#include <cuda/std/cmath>
+
+#include "fp_compare.h"
+#include "test_macros.h"
+
+//  Totally arbitrary picks for precision
+template <typename T>
+__host__ __device__ constexpr T fp_error_pct();
+
+template <>
+__host__ __device__ constexpr float fp_error_pct<float>()
+{
+  return 1.0e-4f;
+}
+
+template <>
+__host__ __device__ constexpr double fp_error_pct<double>()
+{
+  return 1.0e-12;
+}
+
+template <>
+__host__ __device__ constexpr long double fp_error_pct<long double>()
+{
+  return 1.0e-13l;
+}
+
+template <typename T>
+__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<T>::max();
+  constexpr T minV = cuda::std::numeric_limits<T>::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<T>();
+  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<float>();
+  fp_test<double>();
+
+  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
+//
+//===----------------------------------------------------------------------===//
+//
+// <numeric>
+
+// template <class _Tp>
+// _Tp midpoint(_Tp __a, _Tp __b) noexcept
+//
+
+#include <cuda/std/__numeric_>
+#include <cuda/std/cassert>
+#include <cuda/std/cstddef>
+#include <cuda/std/cstdint>
+#include <cuda/std/limits>
+
+#include "test_macros.h"
+
+template <typename T>
+__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<T>;
+
+  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 <typename T>
+__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<T>;
+  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 char>();
+  signed_test<short>();
+  signed_test<int>();
+  signed_test<long>();
+  signed_test<long long>();
+
+  signed_test<cuda::std::int8_t>();
+  signed_test<cuda::std::int16_t>();
+  signed_test<cuda::std::int32_t>();
+  signed_test<cuda::std::int64_t>();
+
+  unsigned_test<unsigned char>();
+  unsigned_test<unsigned short>();
+  unsigned_test<unsigned int>();
+  unsigned_test<unsigned long>();
+  unsigned_test<unsigned long long>();
+
+  unsigned_test<cuda::std::uint8_t>();
+  unsigned_test<cuda::std::uint16_t>();
+  unsigned_test<cuda::std::uint32_t>();
+  unsigned_test<cuda::std::uint64_t>();
+
+#ifndef TEST_HAS_NO_INT128_T
+  unsigned_test<__uint128_t>();
+  signed_test<__int128_t>();
+#endif // !TEST_HAS_NO_INT128_T
+
+  //     int_test<char>();
+  signed_test<cuda::std::ptrdiff_t>();
+  unsigned_test<cuda::std::size_t>();
+
+  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
+//
+//===----------------------------------------------------------------------===//
+//
+// <numeric>
+
+// template <class _Tp>
+// _Tp* midpoint(_Tp* __a, _Tp* __b) noexcept
+//
+
+#include <cuda/std/__numeric_>
+#include <cuda/std/cassert>
+
+#include "test_macros.h"
+
+#if TEST_STD_VER >= 2014
+template <typename T>
+__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 <typename T>
+__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<T>(array, array)), T*);
+  ASSERT_NOEXCEPT(cuda::std::midpoint<T>(array, array));
+  assert(cuda::std::midpoint<T>(array, array) == array);
+  assert(cuda::std::midpoint<T>(array, array + 1000) == array + 500);
+}
+
+template <typename T>
+__host__ __device__ void pointer_test()
+{
+  runtime_test<T>();
+  runtime_test<const T>();
+  runtime_test<volatile T>();
+  runtime_test<const volatile T>();
+
+  //  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<T>();
+  constexpr_test<const T>();
+
+#  if !defined(TEST_COMPILER_GCC)
+  constexpr_test<volatile T>();
+  constexpr_test<const volatile T>();
+#  endif // !TEST_COMPILER_GCC
+#endif // TEST_STD_VER >= 2014 && !TEST_COMPILER_CUDACC_BELOW_11_3
+}
+
+int main(int, char**)
+{
+  pointer_test<char>();
+  pointer_test<int>();
+  pointer_test<double>();
+
+  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
+//
+//===----------------------------------------------------------------------===//
+//
+// <numeric>
+
+// template <class _Tp>
+// _Tp midpoint(_Tp __a, _Tp __b) noexcept
+
+// An overload exists for each of char and all arithmetic types except bool.
+
+#include <cuda/std/__numeric_>
+
+#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 <cuda/std/__algorithm_> // for cuda::std::max
 #include <cuda/std/cassert>
-
-#include <algorithm> // for std::max
-#include <cmath> // for std::abs
+#include <cuda/std/cmath> // 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 <typename T>
-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 <typename T>
-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);
 }