diff --git a/libcudacxx/include/cuda/__numeric/saturating_overflow_cast.h b/libcudacxx/include/cuda/__numeric/saturating_overflow_cast.h
index 94b63864b7a..643c4ebae14 100644
--- a/libcudacxx/include/cuda/__numeric/saturating_overflow_cast.h
+++ b/libcudacxx/include/cuda/__numeric/saturating_overflow_cast.h
@@ -35,13 +35,19 @@ _CCCL_TEMPLATE(class _Up, class _Tp)
 _CCCL_REQUIRES(::cuda::std::__cccl_is_integer_v<_Up> _CCCL_AND ::cuda::std::__cccl_is_integer_v<_Tp>)
 [[nodiscard]] _CCCL_API constexpr overflow_result<_Up> saturating_overflow_cast(_Tp __x) noexcept
 {
-  if (::cuda::std::cmp_less(__x, ::cuda::std::numeric_limits<_Up>::min()))
+  if constexpr (!::cuda::std::in_range<_Up>(::cuda::std::numeric_limits<_Tp>::min()))
   {
-    return {::cuda::std::numeric_limits<_Up>::min(), true};
+    if (::cuda::std::cmp_less(__x, ::cuda::std::numeric_limits<_Up>::min()))
+    {
+      return {::cuda::std::numeric_limits<_Up>::min(), true};
+    }
   }
-  if (::cuda::std::cmp_greater(__x, ::cuda::std::numeric_limits<_Up>::max()))
+  if constexpr (!::cuda::std::in_range<_Up>(::cuda::std::numeric_limits<_Tp>::max()))
   {
-    return {::cuda::std::numeric_limits<_Up>::max(), true};
+    if (::cuda::std::cmp_greater(__x, ::cuda::std::numeric_limits<_Up>::max()))
+    {
+      return {::cuda::std::numeric_limits<_Up>::max(), true};
+    }
   }
   return {static_cast<_Up>(__x), false};
 }
diff --git a/libcudacxx/include/cuda/std/__numeric/saturating_cast.h b/libcudacxx/include/cuda/std/__numeric/saturating_cast.h
index 63098f38c1e..324feb59e68 100644
--- a/libcudacxx/include/cuda/std/__numeric/saturating_cast.h
+++ b/libcudacxx/include/cuda/std/__numeric/saturating_cast.h
@@ -23,15 +23,364 @@
 #include <cuda/__numeric/saturating_overflow_cast.h>
 #include <cuda/std/__concepts/concept_macros.h>
 #include <cuda/std/__type_traits/is_integer.h>
+#include <cuda/std/__type_traits/is_signed.h>
+#include <cuda/std/__type_traits/is_unsigned.h>
+#include <cuda/std/cstdint>
 
 #include <cuda/std/__cccl/prologue.h>
 
 _CCCL_BEGIN_NAMESPACE_CUDA_STD
 
+#if _CCCL_CUDA_COMPILATION()
+_CCCL_TEMPLATE(class _Up, class _Tp)
+_CCCL_REQUIRES((sizeof(_Tp) == sizeof(int8_t)))
+[[nodiscard]] _CCCL_DEVICE_API _Up __saturating_cast_impl_device(_Tp __x, int) noexcept
+{
+  [[maybe_unused]] int __ret;
+
+  if constexpr (sizeof(_Tp) == sizeof(int8_t))
+  {
+    if constexpr (is_signed_v<_Up> && is_unsigned_v<_Tp>)
+    {
+      asm("cvt.sat.s8.u8 %0, %1;" : "=r"(__ret) : "r"(int{__x}));
+      return static_cast<_Up>(__ret);
+    }
+    else if constexpr (is_unsigned_v<_Up> && is_signed_v<_Tp>)
+    {
+      asm("cvt.sat.u8.s8 %0, %1;" : "=r"(__ret) : "r"(int{__x}));
+      return static_cast<_Up>(__ret);
+    }
+    else
+    {
+      return __x;
+    }
+  }
+  else if constexpr (sizeof(_Tp) == sizeof(int16_t))
+  {
+    if constexpr (is_signed_v<_Up> && is_signed_v<_Tp>)
+    {
+      asm("cvt.sat.s8.s16 %0, %1;" : "=r"(__ret) : "h"(__x));
+      return static_cast<_Up>(__ret);
+    }
+    else if constexpr (is_signed_v<_Up> && is_unsigned_v<_Tp>)
+    {
+      asm("cvt.sat.s8.u16 %0, %1;" : "=r"(__ret) : "h"(__x));
+      return static_cast<_Up>(__ret);
+    }
+    else if constexpr (is_unsigned_v<_Up> && is_signed_v<_Tp>)
+    {
+      asm("cvt.sat.u8.s16 %0, %1;" : "=r"(__ret) : "h"(__x));
+      return static_cast<_Up>(__ret);
+    }
+    else
+    {
+      asm("cvt.sat.u8.u16 %0, %1;" : "=r"(__ret) : "h"(__x));
+      return static_cast<_Up>(__ret);
+    }
+  }
+  else if constexpr (sizeof(_Tp) == sizeof(int32_t))
+  {
+    if constexpr (is_signed_v<_Up> && is_signed_v<_Tp>)
+    {
+      asm("cvt.sat.s8.s32 %0, %1;" : "=r"(__ret) : "r"(__x));
+      return static_cast<_Up>(__ret);
+    }
+    else if constexpr (is_signed_v<_Up> && is_unsigned_v<_Tp>)
+    {
+      asm("cvt.sat.s8.u32 %0, %1;" : "=r"(__ret) : "r"(__x));
+      return static_cast<_Up>(__ret);
+    }
+    else if constexpr (is_unsigned_v<_Up> && is_signed_v<_Tp>)
+    {
+      asm("cvt.sat.u8.s32 %0, %1;" : "=r"(__ret) : "r"(__x));
+      return static_cast<_Up>(__ret);
+    }
+    else
+    {
+      asm("cvt.sat.u8.u32 %0, %1;" : "=r"(__ret) : "r"(__x));
+      return static_cast<_Up>(__ret);
+    }
+  }
+  else if constexpr (sizeof(_Tp) == sizeof(int64_t))
+  {
+    if constexpr (is_signed_v<_Up> && is_signed_v<_Tp>)
+    {
+      asm("cvt.sat.s8.s64 %0, %1;" : "=r"(__ret) : "l"(__x));
+      return static_cast<_Up>(__ret);
+    }
+    else if constexpr (is_signed_v<_Up> && is_unsigned_v<_Tp>)
+    {
+      asm("cvt.sat.s8.u64 %0, %1;" : "=r"(__ret) : "l"(__x));
+      return static_cast<_Up>(__ret);
+    }
+    else if constexpr (is_unsigned_v<_Up> && is_signed_v<_Tp>)
+    {
+      asm("cvt.sat.u8.s64 %0, %1;" : "=r"(__ret) : "l"(__x));
+      return static_cast<_Up>(__ret);
+    }
+    else
+    {
+      asm("cvt.sat.u8.u64 %0, %1;" : "=r"(__ret) : "l"(__x));
+      return static_cast<_Up>(__ret);
+    }
+  }
+  else
+  {
+    return ::cuda::saturating_overflow_cast<_Up>(__x).value;
+  }
+}
+
+_CCCL_TEMPLATE(class _Up, class _Tp)
+_CCCL_REQUIRES((sizeof(_Tp) == sizeof(int16_t)))
+[[nodiscard]] _CCCL_DEVICE_API _Up __saturating_cast_impl_device(_Tp __x, int) noexcept
+{
+  [[maybe_unused]] _Up __ret;
+
+  if constexpr (sizeof(_Tp) == sizeof(int8_t))
+  {
+    if constexpr (is_unsigned_v<_Up> && is_signed_v<_Tp>)
+    {
+      asm("cvt.sat.u16.s8 %0, %1;" : "=h"(__ret) : "r"(int{__x}));
+      return __ret;
+    }
+    else
+    {
+      return static_cast<_Up>(__x);
+    }
+  }
+  else if constexpr (sizeof(_Tp) == sizeof(int16_t))
+  {
+    if constexpr (is_signed_v<_Up> && is_unsigned_v<_Tp>)
+    {
+      asm("cvt.sat.s16.u16 %0, %1;" : "=h"(__ret) : "h"(__x));
+      return __ret;
+    }
+    else if constexpr (is_unsigned_v<_Up> && is_signed_v<_Tp>)
+    {
+      asm("cvt.sat.u16.s16 %0, %1;" : "=h"(__ret) : "h"(__x));
+      return __ret;
+    }
+    else
+    {
+      return __x;
+    }
+  }
+  else if constexpr (sizeof(_Tp) == sizeof(int32_t))
+  {
+    if constexpr (is_signed_v<_Up> && is_signed_v<_Tp>)
+    {
+      // There is a bug on Blackwell this PTX instruction giving invalid result for negative inputs. Enable this once
+      // nvbug 6423103 is resolved.
+      NV_IF_ELSE_TARGET(NV_PROVIDES_SM_100, ({ return ::cuda::saturating_overflow_cast<_Up>(__x).value; }), ({
+                          asm("cvt.sat.s16.s32 %0, %1;" : "=h"(__ret) : "r"(__x));
+                          return __ret;
+                        }))
+    }
+    else if constexpr (is_signed_v<_Up> && is_unsigned_v<_Tp>)
+    {
+      asm("cvt.sat.s16.u32 %0, %1;" : "=h"(__ret) : "r"(__x));
+      return __ret;
+    }
+    else if constexpr (is_unsigned_v<_Up> && is_signed_v<_Tp>)
+    {
+      asm("cvt.sat.u16.s32 %0, %1;" : "=h"(__ret) : "r"(__x));
+      return __ret;
+    }
+    else
+    {
+      asm("cvt.sat.u16.u32 %0, %1;" : "=h"(__ret) : "r"(__x));
+      return __ret;
+    }
+  }
+  else if constexpr (sizeof(_Tp) == sizeof(int64_t))
+  {
+    if constexpr (is_signed_v<_Up> && is_signed_v<_Tp>)
+    {
+      // There is a bug on Blackwell this PTX instruction giving invalid result for negative inputs. Enable this once
+      // nvbug 6423103 is resolved.
+      NV_IF_ELSE_TARGET(NV_PROVIDES_SM_100, ({ return ::cuda::saturating_overflow_cast<_Up>(__x).value; }), ({
+                          asm("cvt.sat.s16.s64 %0, %1;" : "=h"(__ret) : "l"(__x));
+                          return __ret;
+                        }))
+    }
+    else if constexpr (is_signed_v<_Up> && is_unsigned_v<_Tp>)
+    {
+      asm("cvt.sat.s16.u64 %0, %1;" : "=h"(__ret) : "l"(__x));
+      return __ret;
+    }
+    else if constexpr (is_unsigned_v<_Up> && is_signed_v<_Tp>)
+    {
+      asm("cvt.sat.u16.s64 %0, %1;" : "=h"(__ret) : "l"(__x));
+      return __ret;
+    }
+    else
+    {
+      asm("cvt.sat.u16.u64 %0, %1;" : "=h"(__ret) : "l"(__x));
+      return __ret;
+    }
+  }
+  else
+  {
+    return ::cuda::saturating_overflow_cast<_Up>(__x).value;
+  }
+}
+
+_CCCL_TEMPLATE(class _Up, class _Tp)
+_CCCL_REQUIRES((sizeof(_Tp) == sizeof(int32_t)))
+[[nodiscard]] _CCCL_DEVICE_API _Up __saturating_cast_impl_device(_Tp __x, int) noexcept
+{
+  [[maybe_unused]] _Up __ret;
+
+  if constexpr (sizeof(_Tp) == sizeof(int8_t))
+  {
+    if constexpr (is_unsigned_v<_Up> && is_signed_v<_Tp>)
+    {
+      asm("cvt.sat.u32.s8 %0, %1;" : "=r"(__ret) : "r"(int{__x}));
+      return __ret;
+    }
+    else
+    {
+      return static_cast<_Up>(__x);
+    }
+  }
+  else if constexpr (sizeof(_Tp) == sizeof(int16_t))
+  {
+    if constexpr (is_unsigned_v<_Up> && is_signed_v<_Tp>)
+    {
+      asm("cvt.sat.u32.s16 %0, %1;" : "=r"(__ret) : "h"(__x));
+      return __ret;
+    }
+    else
+    {
+      return static_cast<_Up>(__x);
+    }
+  }
+  else if constexpr (sizeof(_Tp) == sizeof(int32_t))
+  {
+    if constexpr (is_signed_v<_Up> && is_unsigned_v<_Tp>)
+    {
+      asm("cvt.sat.s32.u32 %0, %1;" : "=r"(__ret) : "r"(__x));
+      return __ret;
+    }
+    else if constexpr (is_unsigned_v<_Up> && is_signed_v<_Tp>)
+    {
+      asm("cvt.sat.u32.s32 %0, %1;" : "=r"(__ret) : "r"(__x));
+      return __ret;
+    }
+    else
+    {
+      return __x;
+    }
+  }
+  else if constexpr (sizeof(_Tp) == sizeof(int64_t))
+  {
+    if constexpr (is_signed_v<_Up> && is_signed_v<_Tp>)
+    {
+      asm("cvt.sat.s32.s64 %0, %1;" : "=r"(__ret) : "l"(__x));
+      return __ret;
+    }
+    else if constexpr (is_signed_v<_Up> && is_unsigned_v<_Tp>)
+    {
+      asm("cvt.sat.s32.u64 %0, %1;" : "=r"(__ret) : "l"(__x));
+      return __ret;
+    }
+    else if constexpr (is_unsigned_v<_Up> && is_signed_v<_Tp>)
+    {
+      asm("cvt.sat.u32.s64 %0, %1;" : "=r"(__ret) : "l"(__x));
+      return __ret;
+    }
+    else
+    {
+      asm("cvt.sat.u32.u64 %0, %1;" : "=r"(__ret) : "l"(__x));
+      return __ret;
+    }
+  }
+  else
+  {
+    return ::cuda::saturating_overflow_cast<_Up>(__x).value;
+  }
+}
+
+_CCCL_TEMPLATE(class _Up, class _Tp)
+_CCCL_REQUIRES((sizeof(_Tp) == sizeof(int64_t)))
+[[nodiscard]] _CCCL_DEVICE_API _Up __saturating_cast_impl_device(_Tp __x, int) noexcept
+{
+  [[maybe_unused]] _Up __ret;
+
+  if constexpr (sizeof(_Tp) == sizeof(int8_t))
+  {
+    if constexpr (is_unsigned_v<_Up> && is_signed_v<_Tp>)
+    {
+      asm("cvt.sat.u64.s8 %0, %1;" : "=l"(__ret) : "r"(int{__x}));
+      return __ret;
+    }
+    else
+    {
+      return static_cast<_Up>(__x);
+    }
+  }
+  else if constexpr (sizeof(_Tp) == sizeof(int16_t))
+  {
+    if constexpr (is_unsigned_v<_Up> && is_signed_v<_Tp>)
+    {
+      asm("cvt.sat.u64.s16 %0, %1;" : "=l"(__ret) : "h"(__x));
+      return __ret;
+    }
+    else
+    {
+      return static_cast<_Up>(__x);
+    }
+  }
+  else if constexpr (sizeof(_Tp) == sizeof(int32_t))
+  {
+    if constexpr (is_unsigned_v<_Up> && is_signed_v<_Tp>)
+    {
+      asm("cvt.sat.u64.s32 %0, %1;" : "=l"(__ret) : "r"(__x));
+      return __ret;
+    }
+    else
+    {
+      return static_cast<_Up>(__x);
+    }
+  }
+  else if constexpr (sizeof(_Tp) == sizeof(int64_t))
+  {
+    if constexpr (is_signed_v<_Up> && is_unsigned_v<_Tp>)
+    {
+      asm("cvt.sat.s64.u64 %0, %1;" : "=l"(__ret) : "l"(__x));
+      return __ret;
+    }
+    else if constexpr (is_unsigned_v<_Up> && is_signed_v<_Tp>)
+    {
+      asm("cvt.sat.u64.s64 %0, %1;" : "=l"(__ret) : "l"(__x));
+      return __ret;
+    }
+    else
+    {
+      return __x;
+    }
+  }
+  else
+  {
+    return ::cuda::saturating_overflow_cast<_Up>(__x).value;
+  }
+}
+
+template <class _Up, class _Tp>
+[[nodiscard]] _CCCL_DEVICE_API _Up __saturating_cast_impl_device(_Tp __x, long) noexcept
+{
+  return ::cuda::saturating_overflow_cast<_Up>(__x).value;
+}
+#endif // _CCCL_CUDA_COMPILATION()
+
 _CCCL_TEMPLATE(class _Up, class _Tp)
 _CCCL_REQUIRES(__cccl_is_integer_v<_Up> _CCCL_AND __cccl_is_integer_v<_Tp>)
 [[nodiscard]] _CCCL_API constexpr _Up saturating_cast(_Tp __x) noexcept
 {
+  _CCCL_IF_NOT_CONSTEVAL_DEFAULT
+  {
+    NV_IF_TARGET(NV_IS_DEVICE, ({ return ::cuda::std::__saturating_cast_impl_device<_Up>(__x, 0); }))
+  }
   return ::cuda::saturating_overflow_cast<_Up>(__x).value;
 }
 
diff --git a/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.sat/saturate_cast.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.sat/saturate_cast.pass.cpp
deleted file mode 100644
index 7a3321a0963..00000000000
--- a/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.sat/saturate_cast.pass.cpp
+++ /dev/null
@@ -1,165 +0,0 @@
-//===----------------------------------------------------------------------===//
-//
-// 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) 2026 NVIDIA CORPORATION & AFFILIATES.
-//
-//===----------------------------------------------------------------------===//
-
-// <cuda/std/numeric>
-
-// template<class R, class T>
-// constexpr R saturating_cast(T x) noexcept;                     // freestanding
-
-#include <cuda/std/cassert>
-#include <cuda/std/limits>
-#include <cuda/std/numeric>
-#include <cuda/std/type_traits>
-
-#include "test_macros.h"
-
-using IMIN = signed char;
-using UMIN = signed char;
-
-#if _CCCL_HAS_INT128()
-using IMAX = __int128_t;
-using UMAX = __uint128_t;
-#else // ^^^ _CCCL_HAS_INT128() ^^^ / vvv !_CCCL_HAS_INT128() vvv
-using IMAX = signed long long;
-using UMAX = unsigned long long;
-#endif // ^^^ !_CCCL_HAS_INT128() ^^^
-
-template <class Ret, class T>
-TEST_FUNC constexpr void test(T x, Ret res, int zero_value)
-{
-  assert(cuda::std::saturating_cast<Ret>(static_cast<T>(zero_value + x)) == res);
-}
-
-template <class S>
-TEST_FUNC constexpr bool test_type(int zero_value)
-{
-  static_assert(cuda::std::is_integral_v<S> && cuda::std::is_signed_v<S>);
-
-  using U = cuda::std::make_unsigned_t<S>;
-
-  constexpr auto small_smax  = cuda::std::numeric_limits<IMIN>::max();
-  constexpr auto small_szero = IMIN{0};
-  constexpr auto small_smin  = cuda::std::numeric_limits<IMIN>::min();
-  constexpr auto small_umax  = cuda::std::numeric_limits<UMIN>::max();
-  constexpr auto small_uzero = UMIN{0};
-
-  constexpr auto big_smax  = cuda::std::numeric_limits<IMAX>::max();
-  constexpr auto big_szero = IMAX{0};
-  constexpr auto big_smin  = cuda::std::numeric_limits<IMAX>::min();
-  constexpr auto big_umax  = cuda::std::numeric_limits<UMAX>::max();
-  constexpr auto big_uzero = UMAX{0};
-
-  constexpr S smax  = cuda::std::numeric_limits<S>::max();
-  constexpr S szero = S{0};
-  constexpr S smin  = cuda::std::numeric_limits<S>::min();
-  constexpr U umax  = cuda::std::numeric_limits<U>::max();
-  constexpr U uzero = U{0};
-
-  // test signed
-
-  static_assert(cuda::std::is_same_v<S, decltype(cuda::std::saturating_cast<S>(small_smax))>);
-  static_assert(noexcept(cuda::std::saturating_cast<S>(small_smax)));
-  test<S>(small_smin, static_cast<S>(small_smin), zero_value);
-  test<S>(small_szero, szero, zero_value);
-  test<S>(small_smax, static_cast<S>(small_smax), zero_value);
-
-  static_assert(cuda::std::is_same_v<S, decltype(cuda::std::saturating_cast<S>(small_umax))>);
-  static_assert(noexcept(cuda::std::saturating_cast<S>(small_umax)));
-  test<S>(small_uzero, szero, zero_value);
-  test<S>(small_umax, (sizeof(S) == sizeof(IMIN)) ? small_smax : static_cast<S>(small_umax), zero_value);
-
-  static_assert(cuda::std::is_same_v<S, decltype(cuda::std::saturating_cast<S>(smax))>);
-  static_assert(noexcept(cuda::std::saturating_cast<S>(smax)));
-  test<S>(smin, smin, zero_value);
-  test<S>(szero, szero, zero_value);
-  test<S>(smax, smax, zero_value);
-
-  static_assert(cuda::std::is_same_v<S, decltype(cuda::std::saturating_cast<S>(umax))>);
-  static_assert(noexcept(cuda::std::saturating_cast<S>(umax)));
-  test<S>(uzero, szero, zero_value);
-  test<S>(umax, smax, zero_value); // saturated
-
-  static_assert(cuda::std::is_same_v<S, decltype(cuda::std::saturating_cast<S>(big_smax))>);
-  static_assert(noexcept(cuda::std::saturating_cast<S>(big_smax)));
-  test<S>(big_smin, smin, zero_value); // saturated
-  test<S>(big_szero, szero, zero_value);
-  test<S>(big_smax, smax, zero_value); // saturated
-
-  static_assert(cuda::std::is_same_v<S, decltype(cuda::std::saturating_cast<S>(big_umax))>);
-  static_assert(noexcept(cuda::std::saturating_cast<S>(big_umax)));
-  test<S>(big_uzero, szero, zero_value);
-  test<S>(big_umax, smax, zero_value); // saturated
-
-  // test unsigned
-
-  static_assert(cuda::std::is_same_v<U, decltype(cuda::std::saturating_cast<U>(small_smax))>);
-  static_assert(noexcept(cuda::std::saturating_cast<U>(small_smax)));
-  test<U>(small_smin, uzero, zero_value);
-  test<U>(small_szero, uzero, zero_value);
-  test<U>(small_smax, static_cast<U>(small_smax), zero_value);
-
-  static_assert(cuda::std::is_same_v<U, decltype(cuda::std::saturating_cast<U>(small_umax))>);
-  static_assert(noexcept(cuda::std::saturating_cast<U>(small_umax)));
-  test<U>(small_uzero, uzero, zero_value);
-  test<U>(small_umax, static_cast<U>(small_umax), zero_value);
-
-  static_assert(cuda::std::is_same_v<U, decltype(cuda::std::saturating_cast<U>(smax))>);
-  static_assert(noexcept(cuda::std::saturating_cast<U>(smax)));
-  test<U>(smin, uzero, zero_value);
-  test<U>(szero, uzero, zero_value);
-  test<U>(smax, static_cast<U>(smax), zero_value);
-
-  static_assert(cuda::std::is_same_v<U, decltype(cuda::std::saturating_cast<U>(umax))>);
-  static_assert(noexcept(cuda::std::saturating_cast<U>(umax)));
-  test<U>(uzero, uzero, zero_value);
-  test<U>(umax, umax, zero_value);
-
-  static_assert(cuda::std::is_same_v<U, decltype(cuda::std::saturating_cast<U>(big_smax))>);
-  static_assert(noexcept(cuda::std::saturating_cast<U>(big_smax)));
-  test<U>(big_smin, uzero, zero_value); // saturated
-  test<U>(big_szero, uzero, zero_value);
-  test<U>(big_smax, (sizeof(U) == sizeof(UMAX)) ? static_cast<U>(smax) : umax, zero_value); // saturated
-
-  static_assert(cuda::std::is_same_v<U, decltype(cuda::std::saturating_cast<U>(big_umax))>);
-  static_assert(noexcept(cuda::std::saturating_cast<U>(big_umax)));
-  test<U>(big_uzero, uzero, zero_value);
-  test<U>(big_umax, umax, zero_value); // saturated
-
-  return true;
-}
-
-TEST_FUNC constexpr bool test(int zero_value)
-{
-  test_type<signed char>(zero_value);
-  test_type<signed short>(zero_value);
-  test_type<signed int>(zero_value);
-  test_type<signed long>(zero_value);
-  test_type<signed long long>(zero_value);
-#if _CCCL_HAS_INT128()
-  test_type<__int128_t>(zero_value);
-#endif // _CCCL_HAS_INT128()
-
-  return true;
-}
-
-__global__ void test_global_kernel(int* zero_value)
-{
-  test(*zero_value);
-  static_assert(test(0));
-}
-
-int main(int, char**)
-{
-  volatile int zero_value = 0;
-
-  test(zero_value);
-  static_assert(test(0));
-
-  return 0;
-}
diff --git a/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.sat/add_sat.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.sat/saturating_add.pass.cpp
similarity index 100%
rename from libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.sat/add_sat.pass.cpp
rename to libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.sat/saturating_add.pass.cpp
diff --git a/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.sat/saturating_cast.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.sat/saturating_cast.pass.cpp
new file mode 100644
index 00000000000..ebd787e0fba
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.sat/saturating_cast.pass.cpp
@@ -0,0 +1,109 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2026 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+// <cuda/std/numeric>
+
+// template<class R, class T>
+// constexpr R saturating_cast(T x) noexcept;                     // freestanding
+
+#include <cuda/std/cassert>
+#include <cuda/std/limits>
+#include <cuda/std/numeric>
+#include <cuda/std/type_traits>
+
+#include "test_macros.h"
+
+template <class To, class From>
+TEST_FUNC constexpr void test(From x, To res, int zero_value)
+{
+  assert(cuda::std::saturating_cast<To>(static_cast<From>(zero_value + x)) == res);
+}
+
+template <class To, class From>
+TEST_FUNC constexpr bool test_type(int zero_value)
+{
+  static_assert(cuda::std::is_same_v<To, decltype(cuda::std::saturating_cast<To>(From{}))>);
+  static_assert(noexcept(cuda::std::saturating_cast<To>(From{})));
+
+  constexpr auto from_min = cuda::std::numeric_limits<From>::min();
+  constexpr auto from_max = cuda::std::numeric_limits<From>::max();
+
+  constexpr auto to_min = cuda::std::numeric_limits<To>::min();
+  constexpr auto to_max = cuda::std::numeric_limits<To>::max();
+
+  test<To>(from_min, (cuda::std::in_range<To>(from_min)) ? static_cast<To>(from_min) : to_min, zero_value);
+  if constexpr (cuda::std::is_signed_v<From>)
+  {
+    test<To>(From{-126}, (cuda::std::is_signed_v<To>) ? static_cast<To>(-126) : To{0}, zero_value);
+    test<To>(From{-1}, (cuda::std::is_signed_v<To>) ? static_cast<To>(-1) : To{0}, zero_value);
+  }
+  test<To>(From{0}, To{0}, zero_value);
+  test<To>(From{1}, To{1}, zero_value);
+  test<To>(From{126}, To{126}, zero_value);
+  test<To>(from_max, (cuda::std::in_range<To>(from_max)) ? static_cast<To>(from_max) : to_max, zero_value);
+
+  return true;
+}
+
+template <class T>
+TEST_FUNC constexpr bool test(int zero_value)
+{
+  test_type<T, signed char>(zero_value);
+  test_type<T, signed short>(zero_value);
+  test_type<T, signed int>(zero_value);
+  test_type<T, signed long>(zero_value);
+  test_type<T, signed long long>(zero_value);
+#if _CCCL_HAS_INT128()
+  test_type<T, __int128_t>(zero_value);
+#endif // _CCCL_HAS_INT128()
+
+  test_type<T, unsigned char>(zero_value);
+  test_type<T, unsigned short>(zero_value);
+  test_type<T, unsigned int>(zero_value);
+  test_type<T, unsigned long>(zero_value);
+  test_type<T, unsigned long long>(zero_value);
+#if _CCCL_HAS_INT128()
+  test_type<T, __uint128_t>(zero_value);
+#endif // _CCCL_HAS_INT128()
+
+  return true;
+}
+
+TEST_FUNC constexpr bool test(int zero_value)
+{
+  test<signed char>(zero_value);
+  test<signed short>(zero_value);
+  test<signed int>(zero_value);
+  test<signed long>(zero_value);
+  test<signed long long>(zero_value);
+#if _CCCL_HAS_INT128()
+  test<__int128_t>(zero_value);
+#endif // _CCCL_HAS_INT128()
+
+  test<unsigned char>(zero_value);
+  test<unsigned short>(zero_value);
+  test<unsigned int>(zero_value);
+  test<unsigned long>(zero_value);
+  test<unsigned long long>(zero_value);
+#if _CCCL_HAS_INT128()
+  test<__uint128_t>(zero_value);
+#endif // _CCCL_HAS_INT128()
+
+  return true;
+}
+
+int main(int, char**)
+{
+  volatile int zero_value = 0;
+
+  test(zero_value);
+  static_assert(test(0));
+
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.sat/div_sat.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.sat/saturating_div.pass.cpp
similarity index 100%
rename from libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.sat/div_sat.pass.cpp
rename to libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.sat/saturating_div.pass.cpp
diff --git a/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.sat/mul_sat.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.sat/saturating_mul.pass.cpp
similarity index 100%
rename from libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.sat/mul_sat.pass.cpp
rename to libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.sat/saturating_mul.pass.cpp
diff --git a/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.sat/sub_sat.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.sat/saturating_sub.pass.cpp
similarity index 100%
rename from libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.sat/sub_sat.pass.cpp
rename to libcudacxx/test/libcudacxx/std/numerics/numeric.ops/numeric.ops.sat/saturating_sub.pass.cpp