[libcudacxx][fp] add fpmp2 multi-precision arithmetic + unit tests - #10517
Draft
akolesov-nvidia wants to merge 34 commits into
Draft
[libcudacxx][fp] add fpmp2 multi-precision arithmetic + unit tests#10517akolesov-nvidia wants to merge 34 commits into
akolesov-nvidia wants to merge 34 commits into
Conversation
Add the fpmp2 class template: a multi-precision floating-point type holding an unevaluated (hi, lo) pair of IEEE-754 components. float elements give the double-float type fp32mp2 (~46 effective mantissa bits) and double elements the double-double type fp64mp2 (~104 bits). An fpmp2_accuracy selector (low / mid / high, def == mid) picks the split and error-accumulation strategy, trading operation count against accuracy; fp32mp2_high reaches ~48 bits. This covers arithmetic only: construction and conversion, comparison, add/sub/mul, div/sqrt, fused multiply-add, atomics, warp shuffle, and numeric_limits. The math functions in <cuda/fpmp_math> follow as a separate PR, so nothing here depends on them. cuda/fpmp public entry point cuda/__fp/fpmp.h fpmp2 class template and type aliases cuda/__fp/fpmp_common.h public accuracy selector and compile-mode knobs cuda/__fp/fpmp_impl*.h arithmetic cores cuda/__fp/fpmp_limits.h numeric_limits specializations 16 lit-style unit tests land under test/libcudacxx/cuda/fp/units/fpmp/.
Contributor
The entry point ended with a trailing blank line, which the end-of-file-fixer pre-commit hook rejects. Matches <cuda/fpemu>, which ends with a single newline.
Caught by the codespell pre-commit hook.
Contributor
Author
|
/ok to test 90a337b |
128-bit atomics: fp64mp2 atomicAdd used atomicCAS(ulonglong2*), an nvcc-only overload that clang-CUDA rejected. It now uses cuda::atomic_ref, guarded by __CUDA_ARCH__ >= 900 && __cccl_ptx_isa >= 840. Targets without a 128-bit compare-exchange no longer get a silent no-op stub: they reference an undefined __fpmp2_dd_atomic_requires_SM_90_and_ptx_isa_840, so such a call fails to link with the requirement named. fp128: ARM64 has no __float128, and _CCCL_FPMP_FP128_ENABLE was keyed on __CUDA_ARCH__ and an arch list rather than on whether the type can be named. Availability now follows _CCCL_HAS_FLOAT128(), and __fpmp_fp128 comes from __fp_native_type_t, so an ARM64 host keeps quad precision through its 128-bit long double -- recognized by format (LDBL_MANT_DIG == 113) instead of by architecture, which also drops the wrong Windows ARM64 case. Consequence: _CCCL_HAS_FLOAT128() is 0 for GCC in strict-ANSI mode, so fp128 is off under nvcc unless CCCL_GCC_HAS_EXTENDED_NUMERIC_LITERALS and -fext-numeric-literals are supplied, matching cuda::std.
Contributor
Author
|
/ok to test 2d3dc3f |
…val build Two fixes on top of the earlier portability work, both from CI. fp128 availability: keying _CCCL_FPMP_FP128_ENABLE on _CCCL_HAS_FLOAT128() was too strict. That macro additionally requires the q/Q literal suffixes, which GCC disables under __STRICT_ANSI__, and on device it requires sm_100+. fpmp needs only the type -- it never writes such literals -- so x86_64 hosts building with -std=c++17 lost quad support they fully have, taking the fp64mp2 math fallback with them. _CCCL_FPMP_HAS_FLOAT128_TYPE now asks the compiler for the type itself, so those builds keep fp128, and a toolchain offering device fp128 below sm_100 can still opt in with -D_CCCL_FPMP_FP128_ENABLE=1. __fpmp_fp128 no longer resolves to void when nothing is available: it names __float128 or a 128-bit long double, and forcing the macro on a platform with neither is an #error stating that, plus a static_assert on the width. _CCCL_FPMP_HOST_SUPPORTS_LDOUBLE128 dropped its _CCCL_HAS_LONG_DOUBLE() term, which is 0 for all of a CUDA compilation including the host pass and so denied ARM64 the only 128-bit type it has; the LDBL_* format test answers the host-ABI question in either pass. ARM64 therefore gets long double, matching what the test suites expect. MSVC: fpmp.h used a local _CCCL_FPMP_IS_CONSTEVAL macro whose non-builtin branch spelled std::is_constant_evaluated(), which inside namespace cuda resolves to cuda::std:: and was never included. Both call sites now use ::cuda::std::is_constant_evaluated(), which is __host__ __device__ and handles the compilers the macro was working around.
Contributor
Author
|
/ok to test 1e11d50 |
This comment has been minimized.
This comment has been minimized.
NVRTC cannot open these headers, so every fpmp test failed to compile in the NVRTC CI job with "catastrophic error: cannot open source file iostream". Both includes were unused: fpmp has no stream-insertion overload, and no fp test, example or benchmark relies on getting them transitively.
Contributor
Author
|
/ok to test 550f770 |
This comment has been minimized.
This comment has been minimized.
MSVC: nvcc's front end reports __cplusplus as 199711L regardless of the /Zc:__cplusplus passed to cl.exe, so the dialect check rejected every MSVC build. Ask _CCCL_STD_VER instead, which reads _MSVC_LANG, and require C++17 outright since the dispatch is written with if constexpr. The same macro now selects the C++20 constructor bodies, which MSVC had silently skipped. MSVC also treats warnings as errors: the round-toward-zero helpers picked nextafterf/nextafter through a ternary, whose common type is double and so narrowed on assignment to float (C4244); if constexpr calls only the matching one. In int.pass.cpp, -2147483648 is unary minus on a literal too large for int (C4146) that then narrows in a braced initializer (C4838). NVRTC: tests reached for the host <type_traits>/<utility> and std:: traits, neither of which exists in a device-only translation unit; they now use the cuda::std equivalents. The three tests that drive kernels through the CUDA runtime API cannot work under NVRTC at all and are marked UNSUPPORTED.
Contributor
Author
|
/ok to test 1f2c3c6 |
This comment has been minimized.
This comment has been minimized.
Use ::cuda::std::isnan instead of std::isnan, which only resolved through a transitive <cmath> include and never through the <cuda/std/cmath> the header actually includes. Drop the unreachable __internal_nv_fpmp2_add_exp variant together with the always-true _CCCL_FPMP_FPAN_METHOD switch that hid it. Correct comments that did not match the code: the documented spelling of the libquadmath override macro, the claim that the warp-shuffle overloads live in the global namespace, and references to fpmp_math/fpmp_lib files and an exp() entry point that are not part of this promotion. Report launch failures in shfl.pass.cpp instead of returning success.
The six fpmp configuration macros date from design exploration and their
defaults are now final, so each knob is removed together with its dead
branch:
* USE_INLINE_ASM_RCP / RSQRT / EX2_LG2: the SFU inline asm becomes
unconditional; the __frcp_rn / __frsqrt_rn / __exp2f / __log2f
fallbacks are gone.
* USE_ACCURATE_DIV was on, so __fpmp2_high_div loses only its guards.
* USE_ACCURATE_MUL was off, so __fpmp2_high_mul and its library-mode
declarations are deleted; mul<high> keeps resolving to the default
path through the existing else arm.
* LARGE_TRIG_FP64_FALLBACK is dropped as an unused definition; the
trig code it gated is not part of this subset.
_CCCL_FPMP_CONSTEXPR and _CCCL_FPMP_NOEXCEPT are inlined as plain
constexpr / noexcept at their few use sites.
Rationale that lived in the deleted knob comments is preserved next to
the code it describes. Behavior is unchanged: host results are
bit-identical before and after.
Contributor
Author
|
/ok to test 9e68be4 |
Follows the fpemu conversion in 16fd3fe: raw `__CUDA_ARCH__` and `__CUDACC__` guards are replaced with the portable CCCL spellings, so host/device selection no longer depends on nvcc's two-pass model. In-function selection becomes `NV_IF_ELSE_TARGET(NV_IS_DEVICE, ...)`, which lets the 35 fp32/fp64 primitives in fpmp_impl.h collapse from a pair of per-pass definitions into one body each. This also fixes nvc++, where `__CUDA_ARCH__` is never defined and device code therefore used the host libm fallbacks instead of the CUDA intrinsics. The 128-bit atomic in fpmp_impl_atomic.h dispatches on `NV_PROVIDES_SM_90`, keeping the named link-time error for older architectures. Declaration-level guards (the warp shuffle overloads, the fp128 enable policy) use `_CCCL_CUDA_COMPILATION()`. Verified bit-identical on host and device (A100, CUDA 12.9) against a pre-change probe over every converted primitive; unit suites pass 31/31 on device and on host.
The __CUDA_ARCH__ conversion routed the double-double atomicAdd through NV_IF_ELSE_TARGET(NV_PROVIDES_SM_90, ...), which broke nvc++: it compiles host and device in one pass, so both arms are instantiated, and cuda::atomic_ref<T> with sizeof(T) == 16 has no valid host instantiation -- its host backend static_asserts on sizeof(T) > 8. Every <cuda/fpmp> TU failed. Availability of a 128-bit compare-exchange therefore has to be decided by the preprocessor here. _CCCL_PTX_ARCH() >= 900 restores the pre-conversion behavior for every compiler (0 outside a device-only pass, so nvc++ takes the named-error path) while keeping raw __CUDA_ARCH__ out of the source, matching how cuda/barrier and cuda/__launch spell arch checks. sm_80 still fails with "Unresolved extern function '__fpmp2_dd_atomic_requires_SM_90_and_ptx_isa_840'" and sm_90 still emits atom.cas.relaxed.gpu.b128. Also converts the six '#ifdef __CUDACC__' declaration guards this file still had to _CCCL_CUDA_COMPILATION(), which the earlier audit's grep missed.
davebayer
reviewed
Jul 30, 2026
This comment has been minimized.
This comment has been minimized.
Spell fpmp2_accuracy::def as mid instead of repeating the numeric value, so the two cannot drift apart. Drop the C++17 #error: libcu++ already refuses to build below its own minimum, and a per-header check only duplicates that. _CCCL_HIDE_FROM_ABI on the defaulted special members, and a mem-initializer list in the volatile copy constructor, following fpemu. Collapse the two C++17/C++20 dialect splits around is_constant_evaluated(). The run-time split is now reached in both dialects: the constructors delegate to a static helper, because through C++17 a constexpr constructor has to initialize every member in its mem-initializer list, so a body that assigns cannot be used in a constant expression there. Delegating keeps one definition and one code path per dialect while staying usable in constexpr context, which the plain body form was not (it broke lut.pass.cpp under C++17). Turn __fpmp_is_fpmp2 into the variable template __fpmp_is_fpmp2_v (35 sites). Add [[nodiscard]] to the pure API: hi/lo, renormalize, the arithmetic and comparison operators, and the free functions. Deliberately not on the atomics or the warp shuffles, whose results are legitimately discarded when the call is made for its side effect, nor on ++/--/compound assignment. Verified: all 18 fpmp tests build and run on the host under C++17/20/23 with -Wall -Wextra -Werror (so a discarded result anywhere in the suite would fail the build), plus device build and run on sm_80 (17 tests; atomic_dd is pre-sm-90) and a device compile of atomic_dd for sm_90.
cuda::ptx::shfl_sync_* is [[nodiscard]] in __ptx/instructions/shfl_sync.h, so the fpmp2 wrappers follow it. The atomics stay discardable, matching cuda::std::atomic::fetch_add and the CUDA atomicAdd they wrap: reading the old value is optional there.
akolesov-nvidia
added a commit
to akolesov-nvidia/cccl
that referenced
this pull request
Jul 30, 2026
Spell fpmp2_accuracy::def as mid instead of repeating the numeric value, so the two cannot drift apart. Drop the C++17 #error: libcu++ already refuses to build below its own minimum, and a per-header check only duplicates that. _CCCL_HIDE_FROM_ABI on the defaulted special members, and a mem-initializer list in the volatile copy constructor, following fpemu. Collapse the two C++17/C++20 dialect splits around is_constant_evaluated(). The run-time split is now reached in both dialects: the constructors delegate to a static helper, because through C++17 a constexpr constructor has to initialize every member in its mem-initializer list, so a body that assigns cannot be used in a constant expression there. Delegating keeps one definition and one code path per dialect while staying usable in constexpr context, which the plain body form was not (it broke lut.pass.cpp under C++17). Turn __fpmp_is_fpmp2 into the variable template __fpmp_is_fpmp2_v (35 sites). Add [[nodiscard]] to the pure API: hi/lo, renormalize, the arithmetic and comparison operators, and the free functions. Deliberately not on the atomics or the warp shuffles, whose results are legitimately discarded when the call is made for its side effect, nor on ++/--/compound assignment. Verified: all 18 fpmp tests build and run on the host under C++17/20/23 with -Wall -Wextra -Werror (so a discarded result anywhere in the suite would fail the build), plus device build and run on sm_80 (17 tests; atomic_dd is pre-sm-90) and a device compile of atomic_dd for sm_90.
The constructor from __fpmp_fp128 was constrained to _FpType == double but the conversion operator was not, so fp32mp2 could convert to a quad yet not be built from one. Under _CCCL_FPMP_USE_LIB that half-conversion was a latent link error: the ABI declares only __fp64mp2_from_quad / __fp64mp2_to_quad, so an fp32 call emitted an undefined reference to __fpmp2_to_quad<float>. Both directions are now fp64mp2 only. fp128 is the interchange type for quad-precision host code (libquadmath, Fortran real(16)) and the pairing that carries meaning is double-double against binary128, 106 significand bits against 113; a double-float holds ~48 bits, fewer than a double, so double is its interchange type. For fp32mp2 the two operations are deleted rather than merely absent, so the diagnostic names the rule: without the deletions the constructor reports an ambiguity (a quad converts to float and double equally well) and the conversion silently routes through operator double(). Callers who want the quad image ask through double, which is exact for any pair within the double-float contract. This mirrors how the 128-bit integer operations are already handled. int.pass.cpp asserts both exclusions alongside the existing fp64mp2 ones. Verified: 20 fpmp tests build and run on the host under C++17/20/23 with -Wall -Wextra -Werror and on sm_80; library mode now references exactly __fp64mp2_from_quad and __fp64mp2_to_quad and no fp32 quad symbol.
Neither precision is a natural fallback for the other, so the component type is always spelled out; the same unused `= float` went from 56 internal primitives.
Reaching the low component by incrementing a component pointer assumed a layout the language does not promise. Library ABI unchanged.
Widening out of fp32mp2 stays implicit, lossy out of fp64mp2 becomes explicit. Also keeps the character constructor's signedness, so char32_t no longer wraps.
It returned the bits of the rounded double, not of the object, and could shadow cuda::std::bit_cast. Use it on hi()/lo() instead; fpemu dropped its own likewise.
Contributor
Author
|
/ok to test ca67e13 |
davebayer
reviewed
Jul 31, 2026
This comment has been minimized.
This comment has been minimized.
Pre-GCC-9 hosts have no __builtin_is_constant_evaluated, so is_constant_evaluated() reads false during constant evaluation and the split walked into the runtime primitive. Default to the cast instead, as cuda::std does. * fpmp: assert each check in the tests instead of accumulating a flag An accumulated ok reported one failure per test with no clue which of its checks broke; assert names the line and the expression, as the fpemu tests already do. * fpmp: say why the kernel launches sit in the NV_IS_HOST branch Read plainly, NV_IF_TARGET(NV_IS_HOST, ...) suggests these tests skip the device. It is the opposite: force_include.h runs main on both, and only the host run can launch, so that branch is the driver of the GPU work. * fpmp: move the atomics' declarations to the atomic header They only sat in fpmp.h so the class could befriend them, and fpmp_impl_atomic.h is already included ahead of the class body. Also records what the constrained conversion templates cost, and drops two doc lines left from the bit_cast removal. (cherry picked from commit 49a32b0) * fpmp: give both compilation passes the same fp128 interface The fp128 gate consulted __CUDA_ARCH__, so nvcc's host pass and its device passes disagreed on whether fpmp2 has a quad constructor and conversion: one definition of the class per pass, and no fp128 interchange at all for host code in a .cu file. Split the question in two. _CCCL_FPMP_FP128_ENABLE now decides only whether the members are declared, from host-toolchain properties that read the same in every pass, and _CCCL_FPMP_FP128_DEVICE_OPS decides whether they carry __device__, folded from __CUDA_ARCH_LIST__ which nvcc exposes in both passes. Below sm_100 they are host functions the device pass parses and discards, so device code asking for quad precision is diagnosed at the call site rather than by an absent member.
Contributor
Author
|
/ok to test be8e990 |
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
GCC 7 rejects a constexpr function whose body mentions a non-constexpr call even when that branch is unreachable, so lut.pass.cpp failed to compile. Select the branch with _CCCL_IF_CONSTEVAL_DEFAULT, which folds to if constexpr when __builtin_is_constant_evaluated is unavailable.
Contributor
Author
|
/ok to test 2d11678 |
This comment has been minimized.
This comment has been minimized.
Contributor
🥳 CI Workflow Results🟩 Finished in 10h 59m: Pass: 100%/115 | Total: 22h 03m | Max: 49m 19s | Hits: 97%/359649See results here. |
davebayer
approved these changes
Aug 3, 2026
Comment on lines
+416
to
+418
| _CCCL_TEMPLATE(typename _Up = _FpType, fpmp2_accuracy _TypeAcc2) | ||
| _CCCL_REQUIRES(__fpmp2_is_fp64_v<_Up>) | ||
| _CCCL_API fpmp2(const fpmp2<float, _TypeAcc2>& __src) noexcept |
Contributor
There was a problem hiding this comment.
Hm, shouldn't this constructor be explicit when _TypeAcc != _TypeAcc2?
Contributor
Author
There was a problem hiding this comment.
yes, I will prepare a fix which generalizes the observation: fp32_lo -> fp64_hi will be explicit, so the tag is opt-in whether the precision changes or not, and fp32_lo -> fp64_lo remains the implicit widening. Thanks for the catch.
… tag A cross-precision conversion accepted any source accuracy tag implicitly, so fp32mp2_low silently became fp64mp2_high while the same-precision cross-accuracy constructor has always been explicit. The tag selects the arithmetic algorithm, so make it opt-in in every direction: widening is implicit only when the tag is preserved, tag-changing conversions are explicit regardless of CCCL_FPMP_EXPLICIT_CASTS, and assignment is offered for the tag-preserving forms only.
Contributor
Author
|
/ok to test eaff1e9 |
Contributor
⏱️ CCCL compile-time benchmark comparison: Public headers compile-time benchResult: 0 regression row(s), 4 improvement row(s) above threshold.
Artifacts: reports and traces TU total compilation
🟢 TU total compilation — Improvements
Direct file processing
🟢 Direct file processing — Improvements
|
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Add the fpmp2 class template: a multi-precision floating-point type holding an unevaluated (hi, lo) pair of IEEE-754 components. float elements give the double-float type fp32mp2 (~46 effective mantissa bits) and double elements the double-double type fp64mp2 (~104 bits). An fpmp2_accuracy selector (low / mid / high, def == mid) picks the split and error-accumulation strategy, trading operation count against accuracy; fp32mp2_high reaches ~48 bits.
This covers arithmetic only: construction and conversion, comparison, add/sub/mul, div/sqrt, fused multiply-add, atomics, warp shuffle, and numeric_limits. The math functions in <cuda/fpmp_math> follow as a separate PR, so nothing here depends on them.
cuda/fpmp public entry point
cuda/__fp/fpmp.h fpmp2 class template and type aliases
cuda/__fp/fpmp_common.h public accuracy selector and compile-mode knobs
cuda/__fp/fpmp_impl*.h arithmetic cores
cuda/__fp/fpmp_limits.h numeric_limits specializations
16 lit-style unit tests land under test/libcudacxx/cuda/fp/units/fpmp/.
Description
closes
Checklist