You signed in with another tab or window. Reload to refresh your session.You signed out in another tab or window. Reload to refresh your session.You switched accounts on another tab or window. Reload to refresh your session.Dismiss alert
No actionable comments were generated in the recent review.
Summary
Summary by CodeRabbit
New Features
Added conditional support for NVIDIA FP8 UE5M3 scalar and two- and four-element vector types, including conversions, classification, numeric limits, and vector and tuple traits. Support requires NVFP8 and CUDA Toolkit 13.4 or later.
Documentation
Documented that UE5M3 is unsupported by DLPack and clarified tuple protocol limitations for its vector types.
Bug Fixes
DLPack conversions now fail at compile time for unsupported UE5M3 values.
Unary minus is unavailable for unsigned CCCL floating-point formats.
This change adds feature-gated NVFP8 UE5M3 scalar and vector support to libcudacxx. It extends floating-point formats, conversions, classification, numeric limits, type traits, and tuple interfaces. It also adds DLPack rejection behavior and expands tests and documentation.
Changes: NVFP8 UE5M3 support
Layer / File(s)
Summary
Feature detection and format representation libcudacxx/include/cuda/std/__cccl/extended_data_types.h, libcudacxx/include/cuda/std/__floating_point/*, libcudacxx/include/cuda/std/__type_traits/*
Adds toolkit-version-gated UE5M3 declarations and a feature macro. Floating-point format properties, storage access, constants, and type traits recognize the format.
Adds UE5M3 conversion paths and tests for casts, conversion ranks, constructors, and operators. A deleted unary-minus overload is added for unsigned __cccl_fp formats.
Adds vector and tuple metadata for UE5M3 x2 and x4 types. DLPack conversion explicitly rejects UE5M3. Documentation describes the DLPack limitation, supported vector types, and tuple protocol restrictions.
Format operations and validation libcudacxx/test/libcudacxx/libcxx/numerics/floating.point/*, libcudacxx/test/libcudacxx/cuda/warp/warp_shuffle.pass.cpp, libcudacxx/test/libcudacxx/libcxx/macros/extended_data_types.pass.cpp
Adds tests for UE5M3 format properties, storage, constants, masks, overflow handling, traits, and warp shuffle.
Suggested reviewers:miscco, fbusato
Change: Feature
Merge Risk: Low · up to 3159f
Some long-double conversions can round incorrectly, and the vector documentation is misleading. These bounded issues merit owner awareness before merging.
Architecture risk: Low · up to 3159f
The change affects 2 systems.
Changed systems:libcudacxx, docs
Architecture concerns
No architecture-level concerns identified.
Review details
Systems and components
observed — libcudacxx (service) was modified; 111 changed files map to changed impact.
observed — docs (service) was modified; 3 changed files map to changed impact.
Before / after behavior
observed — Modified behavior in docs/libcudacxx/extended_api/mdspan/mdspan_to_dlpack.rst: The Constraints section adds __nv_fp8_ue5m3 and its vector types to the documented unsupported types, citing the absence of a corresponding DLPack data type code.
observed — Modified behavior in docs/libcudacxx/extended_api/type_traits/vector_types.rst: The supported extended floating-point vector types list adds __nv_fp8x2_ue5m3 and __nv_fp8x4_ue5m3.
observed — Modified behavior in libcudacxx/include/cuda/__mdspan/mdspan_to_dlpack.h: When _CCCL_HAS_NVFP8_UE5M3() is enabled, __data_type_to_dlpack now matches __nv_fp8_ue5m3 and fails compilation with a dedicated assertion; previously this type had no explicit branch.
observed — Modified behavior in libcudacxx/include/cuda/__type_traits/scalar_type.h: Adds scalar_type specializations for the UE5M3 FP8 two- and four-element vector types under _CCCL_HAS_NVFP8_UE5M3(), each mapping to the scalar ::__nv_fp8_ue5m3.
- [ ] Autofix · Keep fixing CodeRabbit findings and required CI, and resolving merge conflicts
Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 11 remain after this review.
suggestion: Check the value returned by UE5M3 negation.test_fp_neg checks only the return type and noexcept, so this new instantiation cannot detect an incorrect result. Add runtime and static_assert checks for representative UE5M3 values, including a nonzero value. The format is unsigned, and the implementation negates storage by applying the sign mask. (raw.githubusercontent.com) As per coding guidelines, “It is fundamental to verify: ... Runtime and constant-evaluation behavior.”
The cast test only checks 1.0. It does not cover UE5M3 values above the __half range. Add a large UE5M3 input and assert the target formats’ overflow results.
The new source and destination cases both check only 1.f. They cannot detect incorrect handling of zero or values above the __half range. Add UE5M3 boundary cases with independently specified expected results. (raw.githubusercontent.com)
As per coding guidelines, libcudacxx tests must verify “Edge cases” and “Input and output types.”
Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 11 remain after this review.
Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 11 remain after this review.
The reason will be displayed to describe this comment to others. Learn more.
Minor · Functional Correctness · Quick win
important: Avoid double rounding the long double input. For 1.0625L + 0x1p-30L, the intermediate float is 1.0625, so nearest-even conversion produces 1.0 instead of 1.125. The double-source path produces 1.125 for the same value. Preserve the discarded bits when converting to UE5M3, and test values on both sides of a midpoint. (docs.nvidia.com)
As per path instructions: “Focus on correctness” for libcudacxx/**/*.
…DIA#12612NVIDIA#12612 made `cuda::std::fmin` and `cuda::std::fmax` work for all extended
floating-point types by comparing with `cuda::std::isless` and
`cuda::std::isgreater`, and tests them in `fp_min_max.pass.cpp`. Drop our
storage order key and its tests. Values that compare equal, like -0 and +0,
now return the first argument instead of ordering -0 before +0.
Also test `cuda::min`, `cuda::max` and `cuda::clamp` with `__nv_fp8_ue5m3`,
which is supported since NVIDIA#12570.
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
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.
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
CTK 13.4 introduced
__nv_fp8_ue5m3, this PR adds support for this type in libcu++ as we do for similar types.