Enable `__float128` in `cuda::std::ceil`, `floor`, `trunc` and `round` by KassaSana · Pull Request #12253 · NVIDIA/cccl

KassaSana · 2026-10-25T09:51:17Z

Description

closes #9940, closes #9941, closes #9942, closes #9943

Adds __float128 overloads for cuda::std::ceil, floor, trunc and round, following the pattern from #9895 (fabs) and #9991 (copysign):

Host: __builtin_<fn>f128 when available (_CCCL_CHECK_BUILTIN(...) || _CCCL_COMPILER(GCC)), undefined for the nvcc device pass.
Device: __nv_fp128_<fn>.
Fallback: a portable bit-level implementation for hosts that have neither, so every path returns. __trunc_fp128_impl clears the fractional mantissa bits using the existing __fp_* storage and mask helpers. floor, ceil and round are derived from it. round compares x - trunc(x), which is exact, against 0.5 instead of computing trunc(x + 0.5), which avoids double rounding near the halfway point.

The rounding functions are separate overloads rather than one if constexpr template, so the new overloads sit next to the existing __half and __nv_bfloat16 ones in rounding_functions.h.

Tests: the generic test<T>() driver in rounding.pass.cpp also calls llrint, nextafter and the other functions, which don't support __float128 yet. So, like #10014 did in logarithms.pass.cpp, a separate test_float128 driver calls only these four. It also adds float128-specific cases: values beyond the 64-bit range (2^65 + 0.5 and 2^111 + 0.5), halfway cases, signed zeros and NaN.

Verified locally (WSL Ubuntu 24.04, CUDA 13.3; no NVIDIA GPU, so device behavior is up to CI):

rounding.pass.cpp precompiles for sm_100 with a clang 18 host, where _CCCL_HAS_FLOAT128() is 1 on both host and device. It also precompiles with a gcc 13 host, where it is 0 because of __STRICT_ANSI__, as a regression check.
A host-only program compared all four functions bit for bit against libquadmath (ceilq/floorq/truncq/roundq) on special values, every exponent boundary from 2^-4 to 2^116, and 4M random inputs. It was built twice, once with the builtins and once with the builtin defines removed to force the fallback. Both runs had 0 mismatches.
pre-commit run passes on the changed files.

Checklist

New or existing tests cover these changes.
The documentation is up to date with these changes.

Add `__float128` overloads that use `__builtin_<fn>f128` on the host and `__nv_fp128_<fn>` on the device, with a portable bit-level fallback built on a shared truncation helper for hosts that provide neither. The float128 test driver only exercises these four functions, since the other rounding functions do not support `__float128` yet.

copy-pr-bot · 2026-10-25T09:51:20Z

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

Copilot

Copilot was unable to review this pull request because the user who requested the review has insufficient remaining quota.

coderabbitai · 2026-10-25T09:57:49Z

Note

Reviews paused

It looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the reviews.auto_review.auto_pause_after_reviewed_commits setting.

Last reviewed commit: 8c32ca0
Latest commit, not reviewed: 7703cf7

Use the following commands to manage reviews:

@coderabbitai resume to resume automatic reviews.
@coderabbitai review to trigger a single review.

Use the checkboxes below for quick actions:

▶️ Resume reviews
🔍 Trigger review

📝 Summary

Summary by CodeRabbit

New Features
- Added ceil, floor, round, and trunc support for __float128 values when available. These operations preserve fractional precision for large values, handle halfway rounding away from zero, preserve signed-zero results, and propagate NaN values. Availability depends on platform and compilation configuration.

Adds guarded __float128 overloads for ceil, floor, round, and trunc. The implementations use available builtins or device paths, with a portable fallback. Guarded tests cover generic behavior and edge cases for all four functions.

Changes: Float128 rounding support

Layer / File(s)	Summary
Truncation fallback and overload `libcudacxx/include/cuda/std/__cmath/rounding_functions.h`	Adds floating-point format support, `__trunc_fp128_impl`, and the `trunc(__float128)` overload, with builtin, device, and software fallback paths.
Rounding overloads and tests `libcudacxx/include/cuda/std/__cmath/rounding_functions.h`, `libcudacxx/test/libcudacxx/std/numerics/c.math/rounding.pass.cpp`	Adds `ceil(__float128)`, `floor(__float128)`, and `round(__float128)`. Adds guarded tests for all four functions, including large fractional values, halfway rounding, signed zero, and NaN.

Assessment against linked issues

Objective	Addressed	Explanation
Enable `__float128` in `cuda::std::ceil` [9940]	✅
Enable `__float128` in `cuda::std::floor` [9941]	✅
Enable `__float128` in `cuda::std::trunc` [9942]	✅
Enable `__float128` in `cuda::std::round` [9943]	✅

Possibly related issues

NVIDIA/cccl#10471: The PR adds __float128 overloads to cuda::std rounding functions, which relates to extended floating-point support.

Suggested reviewers: miscco

Priority: Low

Change: Feature

Merge Risk: High · up to 8c32c

The new __float128 rounding support references an undeclared variable, so any build with float128 support enabled fails to compile this header. Restoring the one-line storage declaration fixes it, and the build must be fixed before merging.

Finishing Touches: 💡 1

🛠️ Fix failing CI checks: 💡

Commit to this branch
Create a new PR

Autofix · Keep fixing CodeRabbit findings and required CI, and resolving merge conflicts

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

miscco · 2026-10-26T08:34:30Z

+#  if defined(_CCCL_BUILTIN_CEILF128)
+  return _CCCL_BUILTIN_CEILF128(__x);
+#  else // ^^^ _CCCL_BUILTIN_CEILF128 ^^^ / vvv !_CCCL_BUILTIN_CEILF128 vvv
+  NV_IF_TARGET(NV_IS_DEVICE, (return ::__nv_fp128_ceil(__x);))


I believe this requires some NV_PROVIDES_SMXY

Thanks! Switched all four __nv_fp128_* calls to NV_IF_TARGET(NV_PROVIDES_SM_100, ...), matching isnan/min_max, so older architectures fall through to the portable path. Fixed in 65ff009.

miscco

Looks good to me @davebayer can you also give it a look

miscco · 2026-10-26T09:32:09Z

/ok to test 65ff009

coderabbitai

🧹 Nitpick comments: 1

libcudacxx/test/libcudacxx/std/numerics/c.math/rounding.pass.cpp: 1
443-446: Trivial · Functional Correctness · Quick win

suggestion: Cover all four overloads at the 2^111 + 0.5 boundary.

The test checks trunc and ceil at huge, but not floor or round. A boundary regression in either changed overload can pass. Add the corresponding assertions.
   assert(eq(cuda::std::trunc(huge), huge - T(0.5)));
+  assert(eq(cuda::std::floor(huge), huge - T(0.5)));
   assert(eq(cuda::std::ceil(huge), huge + T(0.5)));
+  assert(eq(cuda::std::round(huge), huge + T(0.5)));

Review info

Run configuration

Source mode: Auto
Repository: NVIDIA/cccl
Branch: main
File: .coderabbit.yaml
Commit: 036ecd2
Layers: Repository YAML
Review profile: CHILL
Plan: Enterprise
Run ID: 4c63f65d-23de-4905-9bc8-976d54b51ba7

Commits

Reviewing files that changed from the base of the PR and between 65ff009 and 904103f.

Files selected for processing: 1

libcudacxx/include/cuda/std/__cmath/rounding_functions.h

Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 10 remain after this review.

davebayer · 2026-10-26T12:45:01Z

/ok to test d5dc435

Co-authored-by: David Bayer <[email protected]>

miscco · 2026-10-26T15:11:01Z

/ok to test 8c32ca0

coderabbitai

Actionable comments posted: 1

🧹 Nitpick comments: 1

libcudacxx/include/cuda/std/__cmath/rounding_functions.h: 1
41-41: Trivial · Maintainability & Code Quality · Quick win

suggestion: Mark __trunc_fp128_impl as constexpr.

The helper depends only on pure bit operations and constexpr helpers: __fp_get_exp (decompose.h:63), __fp_get_storage, and __fp_from_storage. The libcudacxx guideline requires constexpr for such functions.
Suggested fix
-[[nodiscard]] _CCCL_API inline __float128 __trunc_fp128_impl(__float128 __x) noexcept
+[[nodiscard]] _CCCL_API constexpr __float128 __trunc_fp128_impl(__float128 __x) noexcept

Review info

Run configuration

Source mode: Auto
Repository: NVIDIA/cccl
Branch: main
File: .coderabbit.yaml
Commit: 036ecd2
Layers: Repository YAML
Review profile: CHILL
Plan: Enterprise
Run ID: 57501574-fa6a-45f0-93f4-04b2381bd7ce

Commits

Reviewing files that changed from the base of the PR and between d5dc435 and 8c32ca0.

Files selected for processing: 1

libcudacxx/include/cuda/std/__cmath/rounding_functions.h

Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 10 remain after this review.

davebayer · 2026-10-26T17:52:25Z

/ok to test 7703cf7

github-actions · 2026-10-26T20:26:55Z

🥳 CI Workflow Results

🟩 Finished in 2h 30m: Pass: 100%/228 | Total: 3d 18h | Max: 1h 30m | Hits: 82%/875804

See results here.

KassaSana requested a review from a team as a code owner October 25, 2026 09:51

Copilot AI balanced review requested due to automatic review settings October 25, 2026 09:51

github-project-automation Bot added this to CCCL Oct 25, 2026

KassaSana requested a review from pciolkosz October 25, 2026 09:51

github-project-automation Bot moved this to Todo in CCCL Oct 25, 2026

Copilot AI reviewed Oct 25, 2026

cccl-authenticator-app Bot moved this from Todo to In Review in CCCL Oct 25, 2026

miscco requested changes Oct 26, 2026

View reviewed changes

github-project-automation Bot moved this from In Review to In Progress in CCCL Oct 26, 2026

Guard __nv_fp128_* rounding intrinsics with NV_PROVIDES_SM_100

65ff009

miscco approved these changes Oct 26, 2026

View reviewed changes

github-project-automation Bot moved this from In Progress to In Review in CCCL Oct 26, 2026

Detect f128 rounding builtins via __has_builtin only, gcc-7 lacks them

904103f

coderabbitai Bot reviewed Oct 26, 2026

View reviewed changes

Cover floor and round at the 2^111 + 0.5 boundary in the float128 test

d5dc435

This comment has been minimized.

Sign in to view

davebayer approved these changes Oct 26, 2026

View reviewed changes

Comment thread libcudacxx/include/cuda/std/__cmath/rounding_functions.h Resolved

This comment has been minimized.

Sign in to view

Use __fp_get_exp

8c32ca0

Co-authored-by: David Bayer <[email protected]>

coderabbitai Bot reviewed Oct 26, 2026

View reviewed changes

Comment thread libcudacxx/include/cuda/std/__cmath/rounding_functions.h Resolved

This comment has been minimized.

Sign in to view

fix

7703cf7

fbusato approved these changes Oct 26, 2026

View reviewed changes

davebayer merged commit ae26e9f into NVIDIA:main Oct 27, 2026
256 checks passed

github-project-automation Bot moved this from In Review to Done in CCCL Oct 27, 2026

This was referenced Oct 28, 2026

[FEA]: Enable __float128 in cuda::std::rint #9947

Closed

Enable __float128 in cuda::std::rint, nearbyint, lrint and llrint #12308

Merged

Enable __float128 in cuda::std::sqrt #12331

Open

coderabbitai Bot mentioned this pull request Oct 29, 2026

[libcu++] Fix __float128 truncation (CI nightly failure) #12339

Merged

This was referenced Nov 3, 2026

Enable __float128 in cuda::std::frexp #12394

Open

Enable __float128 in cuda::std::lround and llround #12398

Merged

Conversation

KassaSana commented Oct 25, 2026

Description

Checklist

Uh oh!

copy-pr-bot Bot commented Oct 25, 2026

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Uh oh!

coderabbitai Bot commented Oct 25, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Reviews paused

Summary by CodeRabbit

Uh oh!

miscco Oct 26, 2026

Choose a reason for hiding this comment

Uh oh!

KassaSana Oct 26, 2026

Choose a reason for hiding this comment

Uh oh!

miscco left a comment

Choose a reason for hiding this comment

Uh oh!

miscco commented Oct 26, 2026

Uh oh!

coderabbitai Bot left a comment

Choose a reason for hiding this comment

Uh oh!

This comment has been minimized.

Uh oh!

davebayer commented Oct 26, 2026

Uh oh!

This comment has been minimized.

miscco commented Oct 26, 2026

Uh oh!

coderabbitai Bot left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

This comment has been minimized.

davebayer commented Oct 26, 2026

Uh oh!

github-actions Bot commented Oct 26, 2026

🥳 CI Workflow Results

🟩 Finished in 2h 30m: Pass: 100%/228 | Total: 3d 18h | Max: 1h 30m | Hits: 82%/875804

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

coderabbitai Bot commented Oct 25, 2026 •

edited

Loading