diff --git a/libcudacxx/include/cuda/std/__cmath/abs.h b/libcudacxx/include/cuda/std/__cmath/abs.h
index 764ebd4eaa4d..fe3e32ffd6c6 100644
--- a/libcudacxx/include/cuda/std/__cmath/abs.h
+++ b/libcudacxx/include/cuda/std/__cmath/abs.h
@@ -87,8 +87,8 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI long double fabsl(long double __x) noe
 template <class _Tp>
 _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __fabs_impl(_Tp __x) noexcept
 {
-  const auto __val = _CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mant_mask_v<_Tp>;
-  return _CUDA_VSTD::__fp_from_storage<_Tp>(static_cast<__fp_storage_t<_Tp>>(__val));
+  const auto __val = _CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mant_mask_of_v<_Tp>;
+  return _CUDA_VSTD::__fp_from_storage<_Tp>(static_cast<__fp_storage_of_t<_Tp>>(__val));
 }
 
 #if _CCCL_HAS_NVFP16()
diff --git a/libcudacxx/include/cuda/std/__cmath/copysign.h b/libcudacxx/include/cuda/std/__cmath/copysign.h
index 552dbafeb914..833b8f81a5b7 100644
--- a/libcudacxx/include/cuda/std/__cmath/copysign.h
+++ b/libcudacxx/include/cuda/std/__cmath/copysign.h
@@ -90,9 +90,9 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __copysign_impl(_Tp __x,
 {
   if constexpr (numeric_limits<_Tp>::is_signed)
   {
-    const auto __val = (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mant_mask_v<_Tp>)
-                     | (_CUDA_VSTD::__fp_get_storage(__y) & __fp_sign_mask_v<_Tp>);
-    return _CUDA_VSTD::__fp_from_storage<_Tp>(static_cast<__fp_storage_t<_Tp>>(__val));
+    const auto __val = (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mant_mask_of_v<_Tp>)
+                     | (_CUDA_VSTD::__fp_get_storage(__y) & __fp_sign_mask_of_v<_Tp>);
+    return _CUDA_VSTD::__fp_from_storage<_Tp>(static_cast<__fp_storage_of_t<_Tp>>(__val));
   }
   else
   {
diff --git a/libcudacxx/include/cuda/std/__cmath/fpclassify.h b/libcudacxx/include/cuda/std/__cmath/fpclassify.h
index 327325efe095..abbe4b1e8bec 100644
--- a/libcudacxx/include/cuda/std/__cmath/fpclassify.h
+++ b/libcudacxx/include/cuda/std/__cmath/fpclassify.h
@@ -86,9 +86,9 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr int __fpclassify_impl(_Tp __
   else
   {
     const auto __storage = _CUDA_VSTD::__fp_get_storage(__x);
-    if ((__storage & __fp_exp_mask_v<_Tp>) == 0)
+    if ((__storage & __fp_exp_mask_of_v<_Tp>) == 0)
     {
-      return (__storage & __fp_mant_mask_v<_Tp>) ? FP_SUBNORMAL : FP_ZERO;
+      return (__storage & __fp_mant_mask_of_v<_Tp>) ? FP_SUBNORMAL : FP_ZERO;
     }
     return FP_NORMAL;
   }
@@ -166,7 +166,7 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr int fpclassify(__nv_fp8_e5m2
 #if _CCCL_HAS_NVFP8_E8M0()
 _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr int fpclassify(__nv_fp8_e8m0 __x) noexcept
 {
-  return ((__x.__x & __fp_exp_mask_v<__nv_fp8_e8m0>) == __fp_exp_mask_v<__nv_fp8_e8m0>) ? FP_NAN : FP_NORMAL;
+  return ((__x.__x & __fp_exp_mask_of_v<__nv_fp8_e8m0>) == __fp_exp_mask_of_v<__nv_fp8_e8m0>) ? FP_NAN : FP_NORMAL;
 }
 #endif // _CCCL_HAS_NVFP8_E8M0()
 
diff --git a/libcudacxx/include/cuda/std/__cmath/isfinite.h b/libcudacxx/include/cuda/std/__cmath/isfinite.h
index 8bd5bd9a2935..6ac3e96039f4 100644
--- a/libcudacxx/include/cuda/std/__cmath/isfinite.h
+++ b/libcudacxx/include/cuda/std/__cmath/isfinite.h
@@ -56,7 +56,7 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isfinite(float __x) noe
     return ::isfinite(__x);
   }
 #  if _LIBCUDACXX_HAS_CONSTEXPR_BIT_CAST()
-  return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mask_v<float>) != __fp_exp_mask_v<float>;
+  return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mask_of_v<float>) != __fp_exp_mask_of_v<float>;
 #  else // ^^^ _LIBCUDACXX_HAS_CONSTEXPR_BIT_CAST() ^^^ / vvv !_LIBCUDACXX_HAS_CONSTEXPR_BIT_CAST() vvv
   return _CUDA_VSTD::__isfinite_impl(__x);
 #  endif // ^^^ !_LIBCUDACXX_HAS_CONSTEXPR_BIT_CAST() ^^^
@@ -73,7 +73,7 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isfinite(double __x) no
     return ::isfinite(__x);
   }
 #  if _LIBCUDACXX_HAS_CONSTEXPR_BIT_CAST()
-  return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mask_v<double>) != __fp_exp_mask_v<double>;
+  return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mask_of_v<double>) != __fp_exp_mask_of_v<double>;
 #  else // ^^^ _LIBCUDACXX_HAS_CONSTEXPR_BIT_CAST() ^^^ / vvv !_LIBCUDACXX_HAS_CONSTEXPR_BIT_CAST() vvv
   return _CUDA_VSTD::__isfinite_impl(__x);
 #  endif // ^^^ !_LIBCUDACXX_HAS_CONSTEXPR_BIT_CAST() ^^^
@@ -94,35 +94,35 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isfinite(long double __
 #if _CCCL_HAS_NVFP16()
 _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isfinite(__half __x) noexcept
 {
-  return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mask_v<__half>) != __fp_exp_mask_v<__half>;
+  return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mask_of_v<__half>) != __fp_exp_mask_of_v<__half>;
 }
 #endif // _CCCL_HAS_NVFP16()
 
 #if _CCCL_HAS_NVBF16()
 _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isfinite(__nv_bfloat16 __x) noexcept
 {
-  return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mask_v<__nv_bfloat16>) != __fp_exp_mask_v<__nv_bfloat16>;
+  return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mask_of_v<__nv_bfloat16>) != __fp_exp_mask_of_v<__nv_bfloat16>;
 }
 #endif // _CCCL_HAS_NVBF16()
 
 #if _CCCL_HAS_NVFP8_E4M3()
 _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isfinite(__nv_fp8_e4m3 __x) noexcept
 {
-  return (__x.__x & __fp_exp_mant_mask_v<__nv_fp8_e4m3>) != __fp_exp_mant_mask_v<__nv_fp8_e4m3>;
+  return (__x.__x & __fp_exp_mant_mask_of_v<__nv_fp8_e4m3>) != __fp_exp_mant_mask_of_v<__nv_fp8_e4m3>;
 }
 #endif // _CCCL_HAS_NVFP8_E4M3()
 
 #if _CCCL_HAS_NVFP8_E5M2()
 _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isfinite(__nv_fp8_e5m2 __x) noexcept
 {
-  return (__x.__x & __fp_exp_mask_v<__nv_fp8_e5m2>) != __fp_exp_mask_v<__nv_fp8_e5m2>;
+  return (__x.__x & __fp_exp_mask_of_v<__nv_fp8_e5m2>) != __fp_exp_mask_of_v<__nv_fp8_e5m2>;
 }
 #endif // _CCCL_HAS_NVFP8_E5M2()
 
 #if _CCCL_HAS_NVFP8_E8M0()
 _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isfinite(__nv_fp8_e8m0 __x) noexcept
 {
-  return __x.__x != __fp_exp_mask_v<__nv_fp8_e8m0>;
+  return __x.__x != __fp_exp_mask_of_v<__nv_fp8_e8m0>;
 }
 #endif // _CCCL_HAS_NVFP8_E8M0()
 
diff --git a/libcudacxx/include/cuda/std/__cmath/isinf.h b/libcudacxx/include/cuda/std/__cmath/isinf.h
index 1c7ff94d0d0c..21f2d62e7893 100644
--- a/libcudacxx/include/cuda/std/__cmath/isinf.h
+++ b/libcudacxx/include/cuda/std/__cmath/isinf.h
@@ -67,7 +67,7 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isinf(float __x) noexce
   {
     return ::isinf(__x);
   }
-  return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mant_mask_v<float>) == __fp_exp_mask_v<float>;
+  return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mant_mask_of_v<float>) == __fp_exp_mask_of_v<float>;
 #else // ^^^ _LIBCUDACXX_HAS_CONSTEXPR_BIT_CAST() ^^^ / vvv !_LIBCUDACXX_HAS_CONSTEXPR_BIT_CAST() vvv
   return _CUDA_VSTD::__isinf_impl(__x);
 #endif // ^^^ !_CCCL_BUILTIN_ISINF ^^^
@@ -89,7 +89,7 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isinf(double __x) noexc
   {
     return ::isinf(__x);
   }
-  return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mant_mask_v<double>) == __fp_exp_mask_v<double>;
+  return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mant_mask_of_v<double>) == __fp_exp_mask_of_v<double>;
 #else // ^^^ _LIBCUDACXX_HAS_CONSTEXPR_BIT_CAST() ^^^ / vvv !_LIBCUDACXX_HAS_CONSTEXPR_BIT_CAST() vvv
   return _CUDA_VSTD::__isinf_impl(__x);
 #endif // ^^^ !_CCCL_BUILTIN_ISINF ^^^
@@ -120,7 +120,7 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isinf(__half __x) noexc
 #    endif // _CCCL_STD_VER <= 2017 || _CCCL_CUDACC_BELOW(12, 3)
   }
 #  endif // _LIBCUDACXX_HAS_NVFP16()
-  return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mant_mask_v<__half>) == __fp_exp_mask_v<__half>;
+  return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mant_mask_of_v<__half>) == __fp_exp_mask_of_v<__half>;
 }
 #endif // _CCCL_HAS_NVFP16()
 
@@ -138,7 +138,8 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isinf(__nv_bfloat16 __x
 #    endif // _CCCL_STD_VER <= 2017 || _CCCL_CUDACC_BELOW(12, 3)
   }
 #  endif // _LIBCUDACXX_HAS_NVBF16()
-  return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mant_mask_v<__nv_bfloat16>) == __fp_exp_mask_v<__nv_bfloat16>;
+  return (_CUDA_VSTD::__fp_get_storage(__x) & __fp_exp_mant_mask_of_v<__nv_bfloat16>)
+      == __fp_exp_mask_of_v<__nv_bfloat16>;
 }
 #endif // _CCCL_HAS_NVBF16()
 
@@ -152,7 +153,7 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isinf(__nv_fp8_e4m3) no
 #if _CCCL_HAS_NVFP8_E5M2()
 _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isinf(__nv_fp8_e5m2 __x) noexcept
 {
-  return (__x.__x & __fp_exp_mant_mask_v<__nv_fp8_e5m2>) == __fp_exp_mask_v<__nv_fp8_e5m2>;
+  return (__x.__x & __fp_exp_mant_mask_of_v<__nv_fp8_e5m2>) == __fp_exp_mask_of_v<__nv_fp8_e5m2>;
 }
 #endif // _CCCL_HAS_NVFP8_E5M2()
 
diff --git a/libcudacxx/include/cuda/std/__cmath/isnan.h b/libcudacxx/include/cuda/std/__cmath/isnan.h
index cd0cdf616dfe..2a928ef7e843 100644
--- a/libcudacxx/include/cuda/std/__cmath/isnan.h
+++ b/libcudacxx/include/cuda/std/__cmath/isnan.h
@@ -85,7 +85,8 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isnan(__half __x) noexc
 #  endif // _LIBCUDACXX_HAS_NVFP16()
 
   const auto __storage = _CUDA_VSTD::__fp_get_storage(__x);
-  return ((__storage & __fp_exp_mask_v<__half>) == __fp_exp_mask_v<__half>) && (__storage & __fp_mant_mask_v<__half>);
+  return ((__storage & __fp_exp_mask_of_v<__half>) == __fp_exp_mask_of_v<__half>)
+      && (__storage & __fp_mant_mask_of_v<__half>);
 }
 #endif // _CCCL_HAS_NVFP16()
 
@@ -100,30 +101,30 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isnan(__nv_bfloat16 __x
 #  endif // _LIBCUDACXX_HAS_NVFP16()
 
   const auto __storage = _CUDA_VSTD::__fp_get_storage(__x);
-  return ((__storage & __fp_exp_mask_v<__nv_bfloat16>) == __fp_exp_mask_v<__nv_bfloat16>)
-      && (__storage & __fp_mant_mask_v<__nv_bfloat16>);
+  return ((__storage & __fp_exp_mask_of_v<__nv_bfloat16>) == __fp_exp_mask_of_v<__nv_bfloat16>)
+      && (__storage & __fp_mant_mask_of_v<__nv_bfloat16>);
 }
 #endif // _CCCL_HAS_NVBF16()
 
 #if _CCCL_HAS_NVFP8_E4M3()
 _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isnan(__nv_fp8_e4m3 __x) noexcept
 {
-  return (__x.__x & __fp_exp_mant_mask_v<__nv_fp8_e4m3>) == __fp_exp_mant_mask_v<__nv_fp8_e4m3>;
+  return (__x.__x & __fp_exp_mant_mask_of_v<__nv_fp8_e4m3>) == __fp_exp_mant_mask_of_v<__nv_fp8_e4m3>;
 }
 #endif // _CCCL_HAS_NVFP8_E4M3()
 
 #if _CCCL_HAS_NVFP8_E5M2()
 _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isnan(__nv_fp8_e5m2 __x) noexcept
 {
-  return ((__x.__x & __fp_exp_mask_v<__nv_fp8_e5m2>) == __fp_exp_mask_v<__nv_fp8_e5m2>)
-      && (__x.__x & __fp_mant_mask_v<__nv_fp8_e5m2>);
+  return ((__x.__x & __fp_exp_mask_of_v<__nv_fp8_e5m2>) == __fp_exp_mask_of_v<__nv_fp8_e5m2>)
+      && (__x.__x & __fp_mant_mask_of_v<__nv_fp8_e5m2>);
 }
 #endif // _CCCL_HAS_NVFP8_E5M2()
 
 #if _CCCL_HAS_NVFP8_E8M0()
 _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isnan(__nv_fp8_e8m0 __x) noexcept
 {
-  return (__x.__x & __fp_exp_mask_v<__nv_fp8_e8m0>) == __fp_exp_mask_v<__nv_fp8_e8m0>;
+  return (__x.__x & __fp_exp_mask_of_v<__nv_fp8_e8m0>) == __fp_exp_mask_of_v<__nv_fp8_e8m0>;
 }
 #endif // _CCCL_HAS_NVFP8_E8M0()
 
diff --git a/libcudacxx/include/cuda/std/__cmath/signbit.h b/libcudacxx/include/cuda/std/__cmath/signbit.h
index 2b35b79a6a77..b03918f2c354 100644
--- a/libcudacxx/include/cuda/std/__cmath/signbit.h
+++ b/libcudacxx/include/cuda/std/__cmath/signbit.h
@@ -68,7 +68,7 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool __signbit_impl(_Tp __x)
 {
   if constexpr (numeric_limits<_Tp>::is_signed)
   {
-    return _CUDA_VSTD::__fp_get_storage(__x) & __fp_sign_mask_v<_Tp>;
+    return _CUDA_VSTD::__fp_get_storage(__x) & __fp_sign_mask_of_v<_Tp>;
   }
   else
   {
diff --git a/libcudacxx/include/cuda/std/__floating_point/cccl_fp.h b/libcudacxx/include/cuda/std/__floating_point/cccl_fp.h
index fad56d827ac9..af1aca0f0008 100644
--- a/libcudacxx/include/cuda/std/__floating_point/cccl_fp.h
+++ b/libcudacxx/include/cuda/std/__floating_point/cccl_fp.h
@@ -35,7 +35,7 @@ class __cccl_fp
 {
   static_assert(_Fmt != __fp_format::__invalid);
 
-  using __storage_type = __fp_storage_t<__cccl_fp>;
+  using __storage_type = __fp_storage_t<_Fmt>;
 
   __storage_type __storage_;
 
@@ -97,9 +97,9 @@ class __cccl_fp
   }
 
   template <class _Tp>
-  friend _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __fp_from_storage(__fp_storage_t<_Tp> __v) noexcept;
+  friend _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __fp_from_storage(__fp_storage_of_t<_Tp> __v) noexcept;
   template <class _Tp>
-  friend _LIBCUDACXX_HIDE_FROM_ABI constexpr __fp_storage_t<_Tp> __fp_get_storage(_Tp __v) noexcept;
+  friend _LIBCUDACXX_HIDE_FROM_ABI constexpr __fp_storage_of_t<_Tp> __fp_get_storage(_Tp __v) noexcept;
 };
 
 _LIBCUDACXX_END_NAMESPACE_STD
diff --git a/libcudacxx/include/cuda/std/__floating_point/mask.h b/libcudacxx/include/cuda/std/__floating_point/mask.h
index 950b020b8f9f..61c3ff212972 100644
--- a/libcudacxx/include/cuda/std/__floating_point/mask.h
+++ b/libcudacxx/include/cuda/std/__floating_point/mask.h
@@ -22,189 +22,38 @@
 #endif // no system header
 
 #include <cuda/std/__floating_point/nvfp_types.h>
+#include <cuda/std/__floating_point/properties.h>
 #include <cuda/std/__floating_point/storage.h>
-#include <cuda/std/__type_traits/always_false.h>
-#include <cuda/std/__type_traits/is_same.h>
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
-template <class _Tp>
-_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr __fp_storage_t<_Tp> __fp_sign_mask_impl() noexcept
-{
-  constexpr __fp_format __fmt = __fp_format_of_v<_Tp>;
-
-  if constexpr (__fmt == __fp_format::__binary16 || __fmt == __fp_format::__bfloat16)
-  {
-    return __fp_storage_t<_Tp>{0x8000u};
-  }
-  else if constexpr (__fmt == __fp_format::__binary32)
-  {
-    return __fp_storage_t<_Tp>{0x80000000u};
-  }
-  else if constexpr (__fmt == __fp_format::__binary64)
-  {
-    return __fp_storage_t<_Tp>{0x8000000000000000ull};
-  }
-  else if constexpr (__fmt == __fp_format::__binary128)
-  {
-    return __fp_storage_t<_Tp>{0x8000000000000000ull} << 64;
-  }
-  else if constexpr (__fmt == __fp_format::__fp80_x86)
-  {
-    return __fp_storage_t<_Tp>{0x8000ull} << 64;
-  }
-  else if constexpr (__fmt == __fp_format::__fp8_nv_e4m3 || __fmt == __fp_format::__fp8_nv_e5m2)
-  {
-    return __fp_storage_t<_Tp>{0x80u};
-  }
-  else if constexpr (__fmt == __fp_format::__fp8_nv_e8m0)
-  {
-    return __fp_storage_t<_Tp>{0x00u};
-  }
-  else if constexpr (__fmt == __fp_format::__fp6_nv_e2m3 || __fmt == __fp_format::__fp6_nv_e3m2)
-  {
-    return __fp_storage_t<_Tp>{0x20u};
-  }
-  else if constexpr (__fmt == __fp_format::__fp4_nv_e2m1)
-  {
-    return __fp_storage_t<_Tp>{0x8u};
-  }
-  else
-  {
-    static_assert(__always_false_v<_Tp>, "Unsupported floating point format");
-  }
-}
-
-template <class _Tp>
-inline constexpr __fp_storage_t<_Tp> __fp_sign_mask_v = __fp_sign_mask_impl<_Tp>();
+template <__fp_format _Fmt>
+inline constexpr auto __fp_sign_mask_v =
+  static_cast<__fp_storage_t<_Fmt>>(__fp_storage_t<_Fmt>(1) << (__fp_exp_nbits_v<_Fmt> + __fp_mant_nbits_v<_Fmt>) );
 
 template <class _Tp>
-_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr __fp_storage_t<_Tp> __fp_exp_mask_impl() noexcept
-{
-  constexpr __fp_format __fmt = __fp_format_of_v<_Tp>;
+inline constexpr auto __fp_sign_mask_of_v = __fp_sign_mask_v<__fp_format_of_v<_Tp>>;
 
-  if constexpr (__fmt == __fp_format::__binary16)
-  {
-    return __fp_storage_t<_Tp>{0x7c00u};
-  }
-  else if constexpr (__fmt == __fp_format::__binary32)
-  {
-    return __fp_storage_t<_Tp>{0x7f800000u};
-  }
-  else if constexpr (__fmt == __fp_format::__binary64)
-  {
-    return __fp_storage_t<_Tp>{0x7ff0000000000000ull};
-  }
-  else if constexpr (__fmt == __fp_format::__binary128)
-  {
-    return __fp_storage_t<_Tp>{0x7fff000000000000ull} << 64;
-  }
-  else if constexpr (__fmt == __fp_format::__bfloat16)
-  {
-    return __fp_storage_t<_Tp>{0x7f80u};
-  }
-  else if constexpr (__fmt == __fp_format::__fp80_x86)
-  {
-    return __fp_storage_t<_Tp>{0x7fffull} << 64;
-  }
-  else if constexpr (__fmt == __fp_format::__fp8_nv_e4m3)
-  {
-    return __fp_storage_t<_Tp>{0x78u};
-  }
-  else if constexpr (__fmt == __fp_format::__fp8_nv_e5m2)
-  {
-    return __fp_storage_t<_Tp>{0x7cu};
-  }
-  else if constexpr (__fmt == __fp_format::__fp8_nv_e8m0)
-  {
-    return __fp_storage_t<_Tp>{0xffu};
-  }
-  else if constexpr (__fmt == __fp_format::__fp6_nv_e2m3)
-  {
-    return __fp_storage_t<_Tp>{0x18u};
-  }
-  else if constexpr (__fmt == __fp_format::__fp6_nv_e3m2)
-  {
-    return __fp_storage_t<_Tp>{0x1cu};
-  }
-  else if constexpr (__fmt == __fp_format::__fp4_nv_e2m1)
-  {
-    return __fp_storage_t<_Tp>{0x6u};
-  }
-  else
-  {
-    static_assert(__always_false_v<_Tp>, "Unsupported floating point format");
-  }
-}
+template <__fp_format _Fmt>
+inline constexpr auto __fp_exp_mask_v = static_cast<__fp_storage_t<_Fmt>>(
+  ((__fp_storage_t<_Fmt>(1) << __fp_exp_nbits_v<_Fmt>) -1) << __fp_mant_nbits_v<_Fmt>);
 
 template <class _Tp>
-inline constexpr __fp_storage_t<_Tp> __fp_exp_mask_v = __fp_exp_mask_impl<_Tp>();
+inline constexpr auto __fp_exp_mask_of_v = __fp_exp_mask_v<__fp_format_of_v<_Tp>>;
 
-template <class _Tp>
-_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr __fp_storage_t<_Tp> __fp_mant_mask_impl() noexcept
-{
-  constexpr __fp_format __fmt = __fp_format_of_v<_Tp>;
-
-  if constexpr (__fmt == __fp_format::__binary16)
-  {
-    return __fp_storage_t<_Tp>{0x03ffu};
-  }
-  else if constexpr (__fmt == __fp_format::__binary32)
-  {
-    return __fp_storage_t<_Tp>{0x007fffffu};
-  }
-  else if constexpr (__fmt == __fp_format::__binary64)
-  {
-    return __fp_storage_t<_Tp>{0x000fffffffffffffull};
-  }
-  else if constexpr (__fmt == __fp_format::__binary128)
-  {
-    return (__fp_storage_t<_Tp>{0x0000ffffffffffffull} << 64) | 0xffffffffffffffffull;
-  }
-  else if constexpr (__fmt == __fp_format::__bfloat16)
-  {
-    return __fp_storage_t<_Tp>{0x007fu};
-  }
-  else if constexpr (__fmt == __fp_format::__fp80_x86)
-  {
-    return __fp_storage_t<_Tp>{0xffffffffffffffffull};
-  }
-  else if constexpr (__fmt == __fp_format::__fp8_nv_e4m3)
-  {
-    return __fp_storage_t<_Tp>{0x07u};
-  }
-  else if constexpr (__fmt == __fp_format::__fp8_nv_e5m2)
-  {
-    return __fp_storage_t<_Tp>{0x03u};
-  }
-  else if constexpr (__fmt == __fp_format::__fp8_nv_e8m0)
-  {
-    return __fp_storage_t<_Tp>{0x00u};
-  }
-  else if constexpr (__fmt == __fp_format::__fp6_nv_e2m3)
-  {
-    return __fp_storage_t<_Tp>{0x07u};
-  }
-  else if constexpr (__fmt == __fp_format::__fp6_nv_e3m2)
-  {
-    return __fp_storage_t<_Tp>{0x03u};
-  }
-  else if constexpr (__fmt == __fp_format::__fp4_nv_e2m1)
-  {
-    return __fp_storage_t<_Tp>{0x01u};
-  }
-  else
-  {
-    static_assert(__always_false_v<_Tp>, "Unsupported floating point format");
-  }
-}
+template <__fp_format _Fmt>
+inline constexpr auto __fp_mant_mask_v =
+  static_cast<__fp_storage_t<_Fmt>>((__fp_storage_t<_Fmt>(1) << __fp_mant_nbits_v<_Fmt>) -1);
 
 template <class _Tp>
-inline constexpr __fp_storage_t<_Tp> __fp_mant_mask_v = __fp_mant_mask_impl<_Tp>();
+inline constexpr auto __fp_mant_mask_of_v = __fp_mant_mask_v<__fp_format_of_v<_Tp>>;
+
+template <__fp_format _Fmt>
+inline constexpr auto __fp_exp_mant_mask_v =
+  static_cast<__fp_storage_t<_Fmt>>(__fp_exp_mask_v<_Fmt> | __fp_mant_mask_v<_Fmt>);
 
 template <class _Tp>
-inline constexpr __fp_storage_t<_Tp> __fp_exp_mant_mask_v =
-  static_cast<__fp_storage_t<_Tp>>(__fp_exp_mask_v<_Tp> | __fp_mant_mask_v<_Tp>);
+inline constexpr auto __fp_exp_mant_mask_of_v = __fp_exp_mant_mask_v<__fp_format_of_v<_Tp>>;
 
 _LIBCUDACXX_END_NAMESPACE_STD
 
diff --git a/libcudacxx/include/cuda/std/__floating_point/native_type.h b/libcudacxx/include/cuda/std/__floating_point/native_type.h
index 782eaaaf7e2e..e2d3eb49093a 100644
--- a/libcudacxx/include/cuda/std/__floating_point/native_type.h
+++ b/libcudacxx/include/cuda/std/__floating_point/native_type.h
@@ -22,7 +22,9 @@
 #endif // no system header
 
 #include <cuda/std/__floating_point/format.h>
+#include <cuda/std/__floating_point/traits.h>
 #include <cuda/std/__type_traits/is_void.h>
+#include <cuda/std/cfloat>
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
@@ -67,6 +69,9 @@ using __fp_native_type_t = decltype(__fp_native_type_impl<_Fmt>());
 template <__fp_format _Fmt>
 inline constexpr bool __fp_has_native_type_v = !_CCCL_TRAIT(is_void, __fp_native_type_t<_Fmt>);
 
+template <class _Tp>
+inline constexpr bool __fp_is_native_type_v = __is_std_fp_v<_Tp> || __is_ext_compiler_fp_v<_Tp>;
+
 _LIBCUDACXX_END_NAMESPACE_STD
 
 #endif // _LIBCUDACXX___FLOATING_POINT_NATIVE_TYPE_H
diff --git a/libcudacxx/include/cuda/std/__floating_point/storage.h b/libcudacxx/include/cuda/std/__floating_point/storage.h
index b90c5a26c740..a4f7f22c3837 100644
--- a/libcudacxx/include/cuda/std/__floating_point/storage.h
+++ b/libcudacxx/include/cuda/std/__floating_point/storage.h
@@ -31,43 +31,44 @@
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
-template <class _Tp>
+template <__fp_format _Fmt>
 _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr auto __fp_storage_type_impl() noexcept
 {
-  constexpr __fp_format __fmt = __fp_format_of_v<_Tp>;
-
-  if constexpr (__fmt == __fp_format::__fp8_nv_e4m3 || __fmt == __fp_format::__fp8_nv_e5m2
-                || __fmt == __fp_format::__fp8_nv_e8m0 || __fmt == __fp_format::__fp6_nv_e2m3
-                || __fmt == __fp_format::__fp6_nv_e3m2 || __fmt == __fp_format::__fp4_nv_e2m1)
+  if constexpr (_Fmt == __fp_format::__fp8_nv_e4m3 || _Fmt == __fp_format::__fp8_nv_e5m2
+                || _Fmt == __fp_format::__fp8_nv_e8m0 || _Fmt == __fp_format::__fp6_nv_e2m3
+                || _Fmt == __fp_format::__fp6_nv_e3m2 || _Fmt == __fp_format::__fp4_nv_e2m1)
   {
     return uint8_t{};
   }
-  else if constexpr (__fmt == __fp_format::__binary16 || __fmt == __fp_format::__bfloat16)
+  else if constexpr (_Fmt == __fp_format::__binary16 || _Fmt == __fp_format::__bfloat16)
   {
     return uint16_t{};
   }
-  else if constexpr (__fmt == __fp_format::__binary32)
+  else if constexpr (_Fmt == __fp_format::__binary32)
   {
     return uint32_t{};
   }
-  else if constexpr (__fmt == __fp_format::__binary64)
+  else if constexpr (_Fmt == __fp_format::__binary64)
   {
     return uint64_t{};
   }
 #if _CCCL_HAS_INT128()
-  else if constexpr (__fmt == __fp_format::__fp80_x86 || __fmt == __fp_format::__binary128)
+  else if constexpr (_Fmt == __fp_format::__fp80_x86 || _Fmt == __fp_format::__binary128)
   {
     return __uint128_t{};
   }
 #endif // _CCCL_HAS_INT128()
   else
   {
-    static_assert(__always_false_v<_Tp>, "Unsupported floating point format");
+    static_assert(__always_false_v<decltype(_Fmt)>, "Unsupported floating point format");
   }
 }
 
+template <__fp_format _Fmt>
+using __fp_storage_t = decltype(__fp_storage_type_impl<_Fmt>());
+
 template <class _Tp>
-using __fp_storage_t = decltype(__fp_storage_type_impl<_Tp>());
+using __fp_storage_of_t = __fp_storage_t<__fp_format_of_v<_Tp>>;
 
 #if _CCCL_HAS_NVFP16()
 struct __cccl_nvfp16_manip_helper : __half
@@ -84,7 +85,7 @@ struct __cccl_nvbf16_manip_helper : __nv_bfloat16
 #endif // _CCCL_HAS_NVBF16()
 
 template <class _Tp>
-_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __fp_from_storage(__fp_storage_t<_Tp> __v) noexcept
+_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __fp_from_storage(__fp_storage_of_t<_Tp> __v) noexcept
 {
   if constexpr (_CCCL_TRAIT(__is_std_fp, _Tp) || _CCCL_TRAIT(__is_ext_compiler_fp, _Tp))
   {
@@ -163,24 +164,22 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __fp_from_storage(__fp_s
     return __ret;
   }
 #endif // _CCCL_HAS_NVFP4_E2M1()
-#if _CCCL_HAS_FLOAT128()
-  else if constexpr (_CCCL_TRAIT(is_same, _Tp, __float128))
-  {
-    return _CUDA_VSTD::bit_cast<_Tp>(__v);
-  }
-#endif // _CCCL_HAS_FLOAT128()
   else
   {
     static_assert(__always_false_v<_Tp>, "Unsupported floating point format");
   }
 }
 
+_CCCL_TEMPLATE(class _Tp, class _Up)
+_CCCL_REQUIRES((!_CCCL_TRAIT(is_same, _Up, __fp_storage_of_t<_Tp>)))
+_LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp __fp_from_storage(const _Up& __v) noexcept = delete;
+
 template <class _Tp>
-_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr __fp_storage_t<_Tp> __fp_get_storage(_Tp __v) noexcept
+_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr __fp_storage_of_t<_Tp> __fp_get_storage(_Tp __v) noexcept
 {
   if constexpr (_CCCL_TRAIT(__is_std_fp, _Tp) || _CCCL_TRAIT(__is_ext_compiler_fp, _Tp))
   {
-    return _CUDA_VSTD::bit_cast<__fp_storage_t<_Tp>>(__v);
+    return _CUDA_VSTD::bit_cast<__fp_storage_of_t<_Tp>>(__v);
   }
   else if constexpr (_CCCL_TRAIT(__is_ext_cccl_fp, _Tp))
   {
@@ -234,12 +233,6 @@ _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr __fp_storage_t<_Tp> __fp_get
     return __v.__x;
   }
 #endif // _CCCL_HAS_NVFP4_E2M1()
-#if _CCCL_HAS_FLOAT128()
-  else if constexpr (_CCCL_TRAIT(is_same, _Tp, __float128))
-  {
-    return _CUDA_VSTD::bit_cast<__fp_storage_t<_Tp>>(__v);
-  }
-#endif // _CCCL_HAS_FLOAT128()
   else
   {
     static_assert(__always_false_v<_Tp>, "Unsupported floating point format");
diff --git a/libcudacxx/include/cuda/std/__limits/numeric_limits_ext.h b/libcudacxx/include/cuda/std/__limits/numeric_limits_ext.h
index 547e87de863d..84522d3ba5c0 100644
--- a/libcudacxx/include/cuda/std/__limits/numeric_limits_ext.h
+++ b/libcudacxx/include/cuda/std/__limits/numeric_limits_ext.h
@@ -45,15 +45,15 @@ class __numeric_limits_impl<__half, __numeric_limits_type::__floating_point>
   static constexpr int max_digits10 = 5;
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type min() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__half>(0x0400u);
+    return _CUDA_VSTD::__fp_from_storage<__half>(uint16_t(0x0400u));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type max() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__half>(0x7bffu);
+    return _CUDA_VSTD::__fp_from_storage<__half>(uint16_t(0x7bffu));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type lowest() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__half>(0xfbffu);
+    return _CUDA_VSTD::__fp_from_storage<__half>(uint16_t(0xfbffu));
   }
 
   static constexpr bool is_integer = false;
@@ -61,11 +61,11 @@ class __numeric_limits_impl<__half, __numeric_limits_type::__floating_point>
   static constexpr int radix       = FLT_RADIX;
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type epsilon() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__half>(0x1400u);
+    return _CUDA_VSTD::__fp_from_storage<__half>(uint16_t(0x1400u));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type round_error() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__half>(0x3800u);
+    return _CUDA_VSTD::__fp_from_storage<__half>(uint16_t(0x3800u));
   }
 
   static constexpr int min_exponent   = -13;
@@ -80,19 +80,19 @@ class __numeric_limits_impl<__half, __numeric_limits_type::__floating_point>
   static constexpr bool has_denorm_loss          = false;
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type infinity() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__half>(0x7c00u);
+    return _CUDA_VSTD::__fp_from_storage<__half>(uint16_t(0x7c00u));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type quiet_NaN() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__half>(0x7e00u);
+    return _CUDA_VSTD::__fp_from_storage<__half>(uint16_t(0x7e00u));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type signaling_NaN() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__half>(0x7d00u);
+    return _CUDA_VSTD::__fp_from_storage<__half>(uint16_t(0x7d00u));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type denorm_min() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__half>(0x0001u);
+    return _CUDA_VSTD::__fp_from_storage<__half>(uint16_t(0x0001u));
   }
 
   static constexpr bool is_iec559  = true;
@@ -122,15 +122,15 @@ class __numeric_limits_impl<__nv_bfloat16, __numeric_limits_type::__floating_poi
   static constexpr int max_digits10 = 4;
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type min() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_bfloat16>(0x0080u);
+    return _CUDA_VSTD::__fp_from_storage<__nv_bfloat16>(uint16_t(0x0080u));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type max() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_bfloat16>(0x7f7fu);
+    return _CUDA_VSTD::__fp_from_storage<__nv_bfloat16>(uint16_t(0x7f7fu));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type lowest() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_bfloat16>(0xff7fu);
+    return _CUDA_VSTD::__fp_from_storage<__nv_bfloat16>(uint16_t(0xff7fu));
   }
 
   static constexpr bool is_integer = false;
@@ -138,11 +138,11 @@ class __numeric_limits_impl<__nv_bfloat16, __numeric_limits_type::__floating_poi
   static constexpr int radix       = FLT_RADIX;
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type epsilon() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_bfloat16>(0x3c00u);
+    return _CUDA_VSTD::__fp_from_storage<__nv_bfloat16>(uint16_t(0x3c00u));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type round_error() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_bfloat16>(0x3f00u);
+    return _CUDA_VSTD::__fp_from_storage<__nv_bfloat16>(uint16_t(0x3f00u));
   }
 
   static constexpr int min_exponent   = -125;
@@ -157,19 +157,19 @@ class __numeric_limits_impl<__nv_bfloat16, __numeric_limits_type::__floating_poi
   static constexpr bool has_denorm_loss          = false;
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type infinity() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_bfloat16>(0x7f80u);
+    return _CUDA_VSTD::__fp_from_storage<__nv_bfloat16>(uint16_t(0x7f80u));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type quiet_NaN() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_bfloat16>(0x7fc0u);
+    return _CUDA_VSTD::__fp_from_storage<__nv_bfloat16>(uint16_t(0x7fc0u));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type signaling_NaN() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_bfloat16>(0x7fa0u);
+    return _CUDA_VSTD::__fp_from_storage<__nv_bfloat16>(uint16_t(0x7fa0u));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type denorm_min() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_bfloat16>(0x0001u);
+    return _CUDA_VSTD::__fp_from_storage<__nv_bfloat16>(uint16_t(0x0001u));
   }
 
   static constexpr bool is_iec559  = true;
@@ -199,15 +199,15 @@ class __numeric_limits_impl<__nv_fp8_e4m3, __numeric_limits_type::__floating_poi
   static constexpr int max_digits10 = 3;
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type min() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e4m3>(0x08u);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e4m3>(uint8_t(0x08u));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type max() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e4m3>(0x7eu);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e4m3>(uint8_t(0x7eu));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type lowest() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e4m3>(0xfeu);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e4m3>(uint8_t(0xfeu));
   }
 
   static constexpr bool is_integer = false;
@@ -215,11 +215,11 @@ class __numeric_limits_impl<__nv_fp8_e4m3, __numeric_limits_type::__floating_poi
   static constexpr int radix       = FLT_RADIX;
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type epsilon() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e4m3>(0x20u);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e4m3>(uint8_t(0x20u));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type round_error() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e4m3>(0x30u);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e4m3>(uint8_t(0x30u));
   }
 
   static constexpr int min_exponent   = -6;
@@ -238,7 +238,7 @@ class __numeric_limits_impl<__nv_fp8_e4m3, __numeric_limits_type::__floating_poi
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type quiet_NaN() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e4m3>(0x7fu);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e4m3>(uint8_t(0x7fu));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type signaling_NaN() noexcept
   {
@@ -246,7 +246,7 @@ class __numeric_limits_impl<__nv_fp8_e4m3, __numeric_limits_type::__floating_poi
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type denorm_min() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e4m3>(0x01u);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e4m3>(uint8_t(0x01u));
   }
 
   static constexpr bool is_iec559  = false;
@@ -276,15 +276,15 @@ class __numeric_limits_impl<__nv_fp8_e5m2, __numeric_limits_type::__floating_poi
   static constexpr int max_digits10 = 2;
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type min() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e5m2>(0x04u);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e5m2>(uint8_t(0x04u));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type max() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e5m2>(0x7bu);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e5m2>(uint8_t(0x7bu));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type lowest() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e5m2>(0xfbu);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e5m2>(uint8_t(0xfbu));
   }
 
   static constexpr bool is_integer = false;
@@ -292,11 +292,11 @@ class __numeric_limits_impl<__nv_fp8_e5m2, __numeric_limits_type::__floating_poi
   static constexpr int radix       = FLT_RADIX;
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type epsilon() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e5m2>(0x34u);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e5m2>(uint8_t(0x34u));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type round_error() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e5m2>(0x38u);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e5m2>(uint8_t(0x38u));
   }
 
   static constexpr int min_exponent   = -15;
@@ -311,19 +311,19 @@ class __numeric_limits_impl<__nv_fp8_e5m2, __numeric_limits_type::__floating_poi
   static constexpr bool has_denorm_loss          = false;
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type infinity() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e5m2>(0x7cu);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e5m2>(uint8_t(0x7cu));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type quiet_NaN() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e5m2>(0x7eu);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e5m2>(uint8_t(0x7eu));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type signaling_NaN() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e5m2>(0x7du);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e5m2>(uint8_t(0x7du));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type denorm_min() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e5m2>(0x01u);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e5m2>(uint8_t(0x01u));
   }
 
   static constexpr bool is_iec559  = false;
@@ -353,15 +353,15 @@ class __numeric_limits_impl<__nv_fp8_e8m0, __numeric_limits_type::__floating_poi
   static constexpr int max_digits10 = 2;
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type min() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e8m0>(0x00u);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e8m0>(uint8_t(0x00u));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type max() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e8m0>(0xfeu);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e8m0>(uint8_t(0xfeu));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type lowest() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e8m0>(0x00u);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e8m0>(uint8_t(0x00u));
   }
 
   static constexpr bool is_integer = false;
@@ -369,11 +369,11 @@ class __numeric_limits_impl<__nv_fp8_e8m0, __numeric_limits_type::__floating_poi
   static constexpr int radix       = FLT_RADIX;
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type epsilon() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e8m0>(0x7fu);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e8m0>(uint8_t(0x7fu));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type round_error() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e8m0>(0x7fu);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e8m0>(uint8_t(0x7fu));
   }
 
   static constexpr int min_exponent   = -127;
@@ -392,7 +392,7 @@ class __numeric_limits_impl<__nv_fp8_e8m0, __numeric_limits_type::__floating_poi
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type quiet_NaN() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e8m0>(0xffu);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp8_e8m0>(uint8_t(0xffu));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type signaling_NaN() noexcept
   {
@@ -430,15 +430,15 @@ class __numeric_limits_impl<__nv_fp6_e2m3, __numeric_limits_type::__floating_poi
   static constexpr int max_digits10 = 3;
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type min() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp6_e2m3>(0x08u);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp6_e2m3>(uint8_t(0x08u));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type max() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp6_e2m3>(0x1fu);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp6_e2m3>(uint8_t(0x1fu));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type lowest() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp6_e2m3>(0x3fu);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp6_e2m3>(uint8_t(0x3fu));
   }
 
   static constexpr bool is_integer = false;
@@ -446,11 +446,11 @@ class __numeric_limits_impl<__nv_fp6_e2m3, __numeric_limits_type::__floating_poi
   static constexpr int radix       = FLT_RADIX;
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type epsilon() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp6_e2m3>(0x01u);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp6_e2m3>(uint8_t(0x01u));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type round_error() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp6_e2m3>(0x04u);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp6_e2m3>(uint8_t(0x04u));
   }
 
   static constexpr int min_exponent   = 0;
@@ -477,7 +477,7 @@ class __numeric_limits_impl<__nv_fp6_e2m3, __numeric_limits_type::__floating_poi
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type denorm_min() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp6_e2m3>(0x01u);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp6_e2m3>(uint8_t(0x01u));
   }
 
   static constexpr bool is_iec559  = false;
@@ -507,15 +507,15 @@ class __numeric_limits_impl<__nv_fp6_e3m2, __numeric_limits_type::__floating_poi
   static constexpr int max_digits10 = 2;
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type min() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp6_e3m2>(0x04u);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp6_e3m2>(uint8_t(0x04u));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type max() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp6_e3m2>(0x1fu);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp6_e3m2>(uint8_t(0x1fu));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type lowest() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp6_e3m2>(0x3fu);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp6_e3m2>(uint8_t(0x3fu));
   }
 
   static constexpr bool is_integer = false;
@@ -523,11 +523,11 @@ class __numeric_limits_impl<__nv_fp6_e3m2, __numeric_limits_type::__floating_poi
   static constexpr int radix       = FLT_RADIX;
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type epsilon() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp6_e3m2>(0x04u);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp6_e3m2>(uint8_t(0x04u));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type round_error() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp6_e3m2>(0x08u);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp6_e3m2>(uint8_t(0x08u));
   }
 
   static constexpr int min_exponent   = -2;
@@ -554,7 +554,7 @@ class __numeric_limits_impl<__nv_fp6_e3m2, __numeric_limits_type::__floating_poi
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type denorm_min() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp6_e3m2>(0x01u);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp6_e3m2>(uint8_t(0x01u));
   }
 
   static constexpr bool is_iec559  = false;
@@ -584,15 +584,15 @@ class __numeric_limits_impl<__nv_fp4_e2m1, __numeric_limits_type::__floating_poi
   static constexpr int max_digits10 = 2;
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type min() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp4_e2m1>(0x2u);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp4_e2m1>(uint8_t(0x2u));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type max() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp4_e2m1>(0x7u);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp4_e2m1>(uint8_t(0x7u));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type lowest() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp4_e2m1>(0xfu);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp4_e2m1>(uint8_t(0xfu));
   }
 
   static constexpr bool is_integer = false;
@@ -600,11 +600,11 @@ class __numeric_limits_impl<__nv_fp4_e2m1, __numeric_limits_type::__floating_poi
   static constexpr int radix       = FLT_RADIX;
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type epsilon() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp4_e2m1>(0x1u);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp4_e2m1>(uint8_t(0x1u));
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type round_error() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp4_e2m1>(0x1u);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp4_e2m1>(uint8_t(0x1u));
   }
 
   static constexpr int min_exponent   = 0;
@@ -631,7 +631,7 @@ class __numeric_limits_impl<__nv_fp4_e2m1, __numeric_limits_type::__floating_poi
   }
   _LIBCUDACXX_HIDE_FROM_ABI static constexpr type denorm_min() noexcept
   {
-    return _CUDA_VSTD::__fp_from_storage<__nv_fp4_e2m1>(0x1u);
+    return _CUDA_VSTD::__fp_from_storage<__nv_fp4_e2m1>(uint8_t(0x1u));
   }
 
   static constexpr bool is_iec559  = false;
diff --git a/libcudacxx/test/libcudacxx/libcxx/numerics/floating.point/mask/exp_mant_mask.pass.cpp b/libcudacxx/test/libcudacxx/libcxx/numerics/floating.point/mask/exp_mant_mask.pass.cpp
new file mode 100644
index 000000000000..1615765cb4cb
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/libcxx/numerics/floating.point/mask/exp_mant_mask.pass.cpp
@@ -0,0 +1,55 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// 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) 2025 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#include <cuda/std/__floating_point/fp.h>
+#include <cuda/std/cassert>
+#include <cuda/std/type_traits>
+
+#include "test_macros.h"
+
+template <cuda::std::__fp_format Fmt>
+__host__ __device__ constexpr bool test_fp_exp_mant_mask(cuda::std::__fp_storage_t<Fmt> expected)
+{
+  using T = cuda::std::__cccl_fp<Fmt>;
+
+  static_assert(cuda::std::is_same_v<cuda::std::remove_cv_t<decltype(cuda::std::__fp_exp_mant_mask_v<Fmt>)>,
+                                     cuda::std::__fp_storage_t<Fmt>>);
+  static_assert(cuda::std::is_same_v<cuda::std::remove_cv_t<decltype(cuda::std::__fp_exp_mant_mask_of_v<T>)>,
+                                     cuda::std::__fp_storage_t<Fmt>>);
+
+  assert(cuda::std::__fp_exp_mant_mask_v<Fmt> == expected);
+  assert(cuda::std::__fp_exp_mant_mask_of_v<T> == expected);
+
+  return true;
+}
+
+static_assert(test_fp_exp_mant_mask<cuda::std::__fp_format::__binary16>(0x7fffu));
+static_assert(test_fp_exp_mant_mask<cuda::std::__fp_format::__binary32>(0x7fffffffu));
+static_assert(test_fp_exp_mant_mask<cuda::std::__fp_format::__binary64>(0x7fffffffffffffffull));
+#if _CCCL_HAS_INT128()
+static_assert(test_fp_exp_mant_mask<cuda::std::__fp_format::__binary128>(
+  (__uint128_t(0x7fffffffffffffffull) << 64) | 0xffffffffffffffffull));
+#endif // _CCCL_HAS_INT128()
+static_assert(test_fp_exp_mant_mask<cuda::std::__fp_format::__bfloat16>(0x7fffu));
+#if _CCCL_HAS_INT128()
+static_assert(
+  test_fp_exp_mant_mask<cuda::std::__fp_format::__fp80_x86>((__uint128_t(0x7fffull) << 64) | 0xffffffffffffffffull));
+#endif // _CCCL_HAS_INT128()
+static_assert(test_fp_exp_mant_mask<cuda::std::__fp_format::__fp8_nv_e4m3>(0x7fu));
+static_assert(test_fp_exp_mant_mask<cuda::std::__fp_format::__fp8_nv_e5m2>(0x7fu));
+static_assert(test_fp_exp_mant_mask<cuda::std::__fp_format::__fp8_nv_e8m0>(0xffu));
+static_assert(test_fp_exp_mant_mask<cuda::std::__fp_format::__fp6_nv_e2m3>(0x1fu));
+static_assert(test_fp_exp_mant_mask<cuda::std::__fp_format::__fp6_nv_e3m2>(0x1fu));
+static_assert(test_fp_exp_mant_mask<cuda::std::__fp_format::__fp4_nv_e2m1>(0x7u));
+
+int main(int, char**)
+{
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/libcxx/numerics/floating.point/mask/exp_mask.pass.cpp b/libcudacxx/test/libcudacxx/libcxx/numerics/floating.point/mask/exp_mask.pass.cpp
new file mode 100644
index 000000000000..8a996b836ac9
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/libcxx/numerics/floating.point/mask/exp_mask.pass.cpp
@@ -0,0 +1,53 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// 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) 2025 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#include <cuda/std/__floating_point/fp.h>
+#include <cuda/std/cassert>
+#include <cuda/std/type_traits>
+
+#include "test_macros.h"
+
+template <cuda::std::__fp_format Fmt>
+__host__ __device__ constexpr bool test_fp_exp_mask(cuda::std::__fp_storage_t<Fmt> expected)
+{
+  using T = cuda::std::__cccl_fp<Fmt>;
+
+  static_assert(cuda::std::is_same_v<cuda::std::remove_cv_t<decltype(cuda::std::__fp_exp_mask_v<Fmt>)>,
+                                     cuda::std::__fp_storage_t<Fmt>>);
+  static_assert(cuda::std::is_same_v<cuda::std::remove_cv_t<decltype(cuda::std::__fp_exp_mask_of_v<T>)>,
+                                     cuda::std::__fp_storage_t<Fmt>>);
+
+  assert(cuda::std::__fp_exp_mask_v<Fmt> == expected);
+  assert(cuda::std::__fp_exp_mask_of_v<T> == expected);
+
+  return true;
+}
+
+static_assert(test_fp_exp_mask<cuda::std::__fp_format::__binary16>(0x7c00u));
+static_assert(test_fp_exp_mask<cuda::std::__fp_format::__binary32>(0x7f800000u));
+static_assert(test_fp_exp_mask<cuda::std::__fp_format::__binary64>(0x7ff0000000000000ull));
+#if _CCCL_HAS_INT128()
+static_assert(test_fp_exp_mask<cuda::std::__fp_format::__binary128>(__uint128_t(0x7fff000000000000ull) << 64));
+#endif // _CCCL_HAS_INT128()
+static_assert(test_fp_exp_mask<cuda::std::__fp_format::__bfloat16>(0x7f80u));
+#if _CCCL_HAS_INT128()
+static_assert(test_fp_exp_mask<cuda::std::__fp_format::__fp80_x86>(__uint128_t(0x7fffull) << 64));
+#endif // _CCCL_HAS_INT128()
+static_assert(test_fp_exp_mask<cuda::std::__fp_format::__fp8_nv_e4m3>(0x78u));
+static_assert(test_fp_exp_mask<cuda::std::__fp_format::__fp8_nv_e5m2>(0x7cu));
+static_assert(test_fp_exp_mask<cuda::std::__fp_format::__fp8_nv_e8m0>(0xffu));
+static_assert(test_fp_exp_mask<cuda::std::__fp_format::__fp6_nv_e2m3>(0x18u));
+static_assert(test_fp_exp_mask<cuda::std::__fp_format::__fp6_nv_e3m2>(0x1cu));
+static_assert(test_fp_exp_mask<cuda::std::__fp_format::__fp4_nv_e2m1>(0x6u));
+
+int main(int, char**)
+{
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/libcxx/numerics/floating.point/mask/mant_mask.pass.cpp b/libcudacxx/test/libcudacxx/libcxx/numerics/floating.point/mask/mant_mask.pass.cpp
new file mode 100644
index 000000000000..8d2fdc23ef97
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/libcxx/numerics/floating.point/mask/mant_mask.pass.cpp
@@ -0,0 +1,54 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// 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) 2025 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#include <cuda/std/__floating_point/fp.h>
+#include <cuda/std/cassert>
+#include <cuda/std/type_traits>
+
+#include "test_macros.h"
+
+template <cuda::std::__fp_format Fmt>
+__host__ __device__ constexpr bool test_fp_mant_mask(cuda::std::__fp_storage_t<Fmt> expected)
+{
+  using T = cuda::std::__cccl_fp<Fmt>;
+
+  static_assert(cuda::std::is_same_v<cuda::std::remove_cv_t<decltype(cuda::std::__fp_mant_mask_v<Fmt>)>,
+                                     cuda::std::__fp_storage_t<Fmt>>);
+  static_assert(cuda::std::is_same_v<cuda::std::remove_cv_t<decltype(cuda::std::__fp_mant_mask_of_v<T>)>,
+                                     cuda::std::__fp_storage_t<Fmt>>);
+
+  assert(cuda::std::__fp_mant_mask_v<Fmt> == expected);
+  assert(cuda::std::__fp_mant_mask_of_v<T> == expected);
+
+  return true;
+}
+
+static_assert(test_fp_mant_mask<cuda::std::__fp_format::__binary16>(0x03ffu));
+static_assert(test_fp_mant_mask<cuda::std::__fp_format::__binary32>(0x007fffffu));
+static_assert(test_fp_mant_mask<cuda::std::__fp_format::__binary64>(0x000fffffffffffffull));
+#if _CCCL_HAS_INT128()
+static_assert(test_fp_mant_mask<cuda::std::__fp_format::__binary128>(
+  (__uint128_t{0x0000ffffffffffffull} << 64) | 0xffffffffffffffffull));
+#endif // _CCCL_HAS_INT128()
+static_assert(test_fp_mant_mask<cuda::std::__fp_format::__bfloat16>(0x007fu));
+#if _CCCL_HAS_INT128()
+static_assert(test_fp_mant_mask<cuda::std::__fp_format::__fp80_x86>(0xffffffffffffffffull));
+#endif // _CCCL_HAS_INT128()
+static_assert(test_fp_mant_mask<cuda::std::__fp_format::__fp8_nv_e4m3>(0x07u));
+static_assert(test_fp_mant_mask<cuda::std::__fp_format::__fp8_nv_e5m2>(0x03u));
+static_assert(test_fp_mant_mask<cuda::std::__fp_format::__fp8_nv_e8m0>(0x00u));
+static_assert(test_fp_mant_mask<cuda::std::__fp_format::__fp6_nv_e2m3>(0x07u));
+static_assert(test_fp_mant_mask<cuda::std::__fp_format::__fp6_nv_e3m2>(0x03u));
+static_assert(test_fp_mant_mask<cuda::std::__fp_format::__fp4_nv_e2m1>(0x1u));
+
+int main(int, char**)
+{
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/libcxx/numerics/floating.point/mask/sign_mask.pass.cpp b/libcudacxx/test/libcudacxx/libcxx/numerics/floating.point/mask/sign_mask.pass.cpp
new file mode 100644
index 000000000000..77dc2685a4fd
--- /dev/null
+++ b/libcudacxx/test/libcudacxx/libcxx/numerics/floating.point/mask/sign_mask.pass.cpp
@@ -0,0 +1,53 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// 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) 2025 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#include <cuda/std/__floating_point/fp.h>
+#include <cuda/std/cassert>
+#include <cuda/std/type_traits>
+
+#include "test_macros.h"
+
+template <cuda::std::__fp_format Fmt>
+__host__ __device__ constexpr bool test_fp_sign_mask(cuda::std::__fp_storage_t<Fmt> expected)
+{
+  using T = cuda::std::__cccl_fp<Fmt>;
+
+  static_assert(cuda::std::is_same_v<cuda::std::remove_cv_t<decltype(cuda::std::__fp_sign_mask_v<Fmt>)>,
+                                     cuda::std::__fp_storage_t<Fmt>>);
+  static_assert(cuda::std::is_same_v<cuda::std::remove_cv_t<decltype(cuda::std::__fp_sign_mask_of_v<T>)>,
+                                     cuda::std::__fp_storage_t<Fmt>>);
+
+  assert(cuda::std::__fp_sign_mask_v<Fmt> == expected);
+  assert(cuda::std::__fp_sign_mask_of_v<T> == expected);
+
+  return true;
+}
+
+static_assert(test_fp_sign_mask<cuda::std::__fp_format::__binary16>(0x8000u));
+static_assert(test_fp_sign_mask<cuda::std::__fp_format::__binary32>(0x80000000u));
+static_assert(test_fp_sign_mask<cuda::std::__fp_format::__binary64>(0x8000000000000000ull));
+#if _CCCL_HAS_INT128()
+static_assert(test_fp_sign_mask<cuda::std::__fp_format::__binary128>(__uint128_t(0x8000000000000000ull) << 64));
+#endif // _CCCL_HAS_INT128()
+static_assert(test_fp_sign_mask<cuda::std::__fp_format::__bfloat16>(0x8000u));
+#if _CCCL_HAS_INT128()
+static_assert(test_fp_sign_mask<cuda::std::__fp_format::__fp80_x86>(__uint128_t(0x8000ull) << 64));
+#endif // _CCCL_HAS_INT128()
+static_assert(test_fp_sign_mask<cuda::std::__fp_format::__fp8_nv_e4m3>(0x80u));
+static_assert(test_fp_sign_mask<cuda::std::__fp_format::__fp8_nv_e5m2>(0x80u));
+static_assert(test_fp_sign_mask<cuda::std::__fp_format::__fp8_nv_e8m0>(0x00u));
+static_assert(test_fp_sign_mask<cuda::std::__fp_format::__fp6_nv_e2m3>(0x20u));
+static_assert(test_fp_sign_mask<cuda::std::__fp_format::__fp6_nv_e3m2>(0x20u));
+static_assert(test_fp_sign_mask<cuda::std::__fp_format::__fp4_nv_e2m1>(0x8u));
+
+int main(int, char**)
+{
+  return 0;
+}
diff --git a/libcudacxx/test/libcudacxx/libcxx/numerics/floating.point/storage.pass.cpp b/libcudacxx/test/libcudacxx/libcxx/numerics/floating.point/storage.pass.cpp
index 035c450994e7..84e879a52903 100644
--- a/libcudacxx/test/libcudacxx/libcxx/numerics/floating.point/storage.pass.cpp
+++ b/libcudacxx/test/libcudacxx/libcxx/numerics/floating.point/storage.pass.cpp
@@ -23,7 +23,7 @@
 template <class T>
 __host__ __device__ void test_fp_storage()
 {
-  using Storage = cuda::std::__fp_storage_t<T>;
+  using Storage = cuda::std::__fp_storage_of_t<T>;
   static_assert(cuda::std::is_integral_v<Storage>);
   static_assert(sizeof(Storage) == sizeof(T));
   static_assert(alignof(Storage) == alignof(T));