[libcu++] Add __float128 support for cuda::std::fabs

[davebayer]

Fixes #9903

[coderabbitai]

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: f2540c3d-3af5-4c99-9a6b-69a3d15350e8

📥 Commits

Reviewing files that changed from the base of the PR and between 895d465 and a9199a1.

📒 Files selected for processing (4)

• libcudacxx/include/cuda/std/__cmath/abs.h
• libcudacxx/include/cuda/std/__floating_point/cast.h
• libcudacxx/test/libcudacxx/libcxx/numerics/floating.point/cast.pass.cpp
• libcudacxx/test/libcudacxx/std/numerics/c.math/fp_abs.pass.cpp

🚧 Files skipped from review as they are similar to previous changes (3)

• libcudacxx/test/libcudacxx/std/numerics/c.math/fp_abs.pass.cpp
• libcudacxx/include/cuda/std/__cmath/abs.h
• libcudacxx/include/cuda/std/__floating_point/cast.h

---

📝 Walkthrough

Summary by CodeRabbit

• New Features
  • Added/extended __float128 support for absolute-value calculations when available.
  • Improved floating-point casting to and from __float128, including conversions across common floating formats.
• Tests
  • Expanded __float128 coverage for floating-point casting and fabs/abs behavior when enabled.
  • Updated host-side test compilation to support extended numeric literals.

Walkthrough

Changes

Adds guarded __float128 support to cuda::std::fabs and __fp_cast, including compiler-specific fallback handling. Tests cover __float128 casts and absolute-value operations when supported.

__float128 support

Layer / File(s)	Summary
__float128 fabs handling libcudacxx/include/cuda/std/__cmath/abs.h	Adds guarded builtin and CUDA fallback paths for fabs(__float128), including constexpr handling.
__float128 conversion dispatch libcudacxx/include/cuda/std/__floating_point/cast.h	Adds conversions to and from __float128 across supported floating-point formats.
__float128 test coverage libcudacxx/test/libcudacxx/.../cast.pass.cpp, libcudacxx/test/libcudacxx/.../fp_abs.pass.cpp	Enables extended numeric literals and adds feature-guarded cast and absolute-value tests.

Suggested labels: libcu++

Suggested reviewers: pciolkosz, miscco

---

_{Comment @coderabbitai help to get the list of available commands.}

[coderabbitai]

Actionable comments posted: 1

🧹 Nitpick comments (1)

libcudacxx/include/cuda/std/__floating_point/cast.h (1)

868-948: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

suggestion: __float128→half/bf16/fp8/fp6/fp4 branches duplicate the float-source dispatch table instead of delegating.

Every other non-float source branch in this function (e.g. __half, __nv_bfloat16, the fp8/fp6/fp4 sources) reaches these target types via ::cuda::std::__fp_cast<_To>(::cuda::std::__fp_cast<float>(__v)), delegating to the existing float-source dispatch. The new __float128 branch instead re-implements each target's raw intrinsic call (::__float2half, ::__float2bfloat16, ::__nv_cvt_float_to_fp8, etc.) directly on static_cast<float>(__v), duplicating logic that already lives in the float-source branch above.

Suggested simplification

-#  if _CCCL_HAS_NVFP16()
-    else if constexpr (is_same_v<_To, __half>)
-    {
-      return ::__float2half(static_cast<float>(__v));
-    }
-#  endif // _CCCL_HAS_NVFP16()
-#  if _CCCL_HAS_NVBF16()
-    else if constexpr (is_same_v<_To, __nv_bfloat16>)
-    {
-      return ::__float2bfloat16(static_cast<float>(__v));
-    }
-#  endif // _CCCL_HAS_NVBF16()
-#  if _CCCL_HAS_NVFP8_E4M3()
-    else if constexpr (is_same_v<_To, __nv_fp8_e4m3>)
-    {
-      return ::cuda::std::__fp_from_storage<__nv_fp8_e4m3>(
-        ::__nv_cvt_float_to_fp8(static_cast<float>(__v), __NV_NOSAT, __NV_E4M3));
-    }
-#  endif // _CCCL_HAS_NVFP8_E4M3()
-    ... (fp8_e5m2, fp8_e8m0, fp6_e2m3, fp6_e3m2, fp4_e2m1 likewise)
+    else if constexpr (is_same_v<_To, __float128>)
+    {
+      return __v;
+    }
+    else
+    {
+      return ::cuda::std::__fp_cast<_To>(static_cast<float>(__v));
+    }

---

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: cc8d5b4c-51a7-4a37-bc84-f3e137d37303

📥 Commits

Reviewing files that changed from the base of the PR and between 0d919ed and b886cfb.

📒 Files selected for processing (4)

• libcudacxx/include/cuda/std/__cmath/abs.h
• libcudacxx/include/cuda/std/__floating_point/cast.h
• libcudacxx/test/libcudacxx/libcxx/numerics/floating.point/cast.pass.cpp
• libcudacxx/test/libcudacxx/std/numerics/c.math/fp_abs.pass.cpp

[coderabbitai]

Actionable comments posted: 1

---

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: aec86913-01e8-4ce7-a021-277d2f9d2421

📥 Commits

Reviewing files that changed from the base of the PR and between b886cfb and 895d465.

📒 Files selected for processing (4)

• libcudacxx/include/cuda/std/__cmath/abs.h
• libcudacxx/include/cuda/std/__floating_point/cast.h
• libcudacxx/test/libcudacxx/libcxx/numerics/floating.point/cast.pass.cpp
• libcudacxx/test/libcudacxx/std/numerics/c.math/fp_abs.pass.cpp

🚧 Files skipped from review as they are similar to previous changes (3)

• libcudacxx/test/libcudacxx/libcxx/numerics/floating.point/cast.pass.cpp
• libcudacxx/test/libcudacxx/std/numerics/c.math/fp_abs.pass.cpp
• libcudacxx/include/cuda/std/__floating_point/cast.h

[github-actions]

🥳 CI Workflow Results

🟩 Finished in 2h 11m: Pass: 100%/120 | Total: 3d 15h | Max: 2h 06m | Hits: 43%/1154054

See results here.

[fbusato]

approved. very minor concern if the future the builtin will be supported on device code, but the SASS differs from the intrinsic. (I would prioritize the intrinsic)