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
{{ message }}
Repository navigation
[libcu++] Add cuda::min, cuda::max and cuda::clamp - #12512
cuda::std::min, cuda::std::max and cuda::std::clamp must take their arguments by reference and return a reference to one of them, so we cannot change them the way the first version of this PR did. Instead, this adds cuda::min, cuda::max and cuda::clamp to <cuda/cmath>, which take and return values. The compiler can then compute the result in registers instead of selecting the address of an argument and loading the result from it, so on device cuda::clamp maps to a single min and max instruction. The exception is double on GPUs with slow FP64 arithmetic, where cuda::std::fmin and cuda::std::fmax compare with integer instructions since #12231.
The functions accept arithmetic types and extended floating-point types, including the FP8, FP6 and FP4 types that cuda::std::fmin and cuda::std::fmax support since [libcu++] Make fmin/fmax work for all nvfp types #12612. All arguments must have the same type.
Floating-point values are compared like with cuda::std::fmin and cuda::std::fmax, so a NaN argument is ignored. cuda::clamp(v, lo, hi) is cuda::max(cuda::min(v, hi), lo): a NaN bound is ignored, and a NaN v is clamped to hi, or to lo if hi is NaN.
cuda::min and cuda::max call cuda::minimum and cuda::maximum, so there is a single implementation. Because they take their arguments by value, the references that cuda::minimum and cuda::maximum take are optimized away: the SASS for float, double, __half and __nv_bfloat16 is the same as when calling cuda::std::fmin and cuda::std::fmax directly.
Through cuda::minimum and cuda::maximum, the new functions use cuda::std::fmin and cuda::std::fmax, which this also improves:
The functions of <cuda_fp16.h> and <cuda_bf16.h> are called unqualified and found by ADL, so <cuda/std/__cmath/min_max.h> does not need to include those headers.
Before SM80, __nv_bfloat16 uses min.f32 and max.f32 instead of the emulation in __hmin and __hmax, which takes about 50 instructions with branches. The result is exact, which was verified for all 2^32 input pairs.
Before CTK 12.2, __hmin and __hmax are only available on device, so the host compares as float.
Tests and documentation:
New cuda/cmath/min_max.pass.cpp and cuda/cmath/clamp.pass.cpp, which run with the conversions and operators of the extended floating-point types disabled (__CUDA_NO_HALF_OPERATORS__ and friends).
A new cmath-ptx FileCheck suite in libcudacxx/test/cmath_codegen, which checks that cuda::min, cuda::max and cuda::clamp compile to min and max instructions without branches. For double, the instructions are only checked on SM80, SM90 and SM100, because the other targets in CI have slow FP64 arithmetic. It is wired into CI like simd-ptx.
New documentation pages for cuda::min and cuda::max, and for cuda::clamp.
Checklist
New or existing tests cover these changes.
The documentation is up to date with these changes.
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.
No actionable comments were generated in the recent review.
Summary
Summary by CodeRabbit
New Features
Added cuda::clamp, cuda::min, and cuda::max for supported arithmetic types, including extended floating-point types.
Floating-point min/max handle NaN inputs. cuda::clamp documents NaN behavior and validates its bounds.
Updated fmin and fmax behavior for NVIDIA extended floating-point formats, including FP8, FP6, and FP4.
Documentation
Added reference pages and examples for the math functions, with availability information for CCCL 3.7.0 and CUDA 13.7.
The PR adds public cuda::min, cuda::max, and cuda::clamp APIs. It updates extended floating-point fmin and fmax paths, adds functional and PTX codegen coverage, documents the APIs, and enables the cmath PTX suite in presets and CI matrices.
Changes: CUDA math APIs
Layer / File(s)
Summary
Min/max and clamp APIs libcudacxx/include/cuda/std/__cmath/min_max.h, libcudacxx/include/cuda/__cmath/min_max.h, libcudacxx/include/cuda/__cmath/clamp.h, libcudacxx/include/cuda/cmath, docs/libcudacxx/extended_api/math.rst, docs/libcudacxx/extended_api/math/*
The changes add cuda::min, cuda::max, and cuda::clamp, export their headers, and document their behavior and availability. The extended floating-point fmin and fmax paths add storage-order comparisons and target-specific handling for half and bfloat16.
API and floating-point tests libcudacxx/test/libcudacxx/cuda/cmath/*, libcudacxx/test/libcudacxx/std/numerics/c.math/fp_min_max.pass.cpp
Tests check API constraints, result types, noexcept, and results across built-in and supported extended types. Narrow-format tests compare fmin and fmax results by storage representation.
PTX codegen coverage and CI wiring libcudacxx/test/cmath_codegen/*, libcudacxx/test/CMakeLists.txt, CMakePresets.json, ci/matrix.yaml
The new PTX tests check instruction patterns for clamp and min/max. CMake registers the cmath codegen suite, and presets and CI matrices include its target.
Assessment against linked issues
Objective
Addressed
Explanation
Provide a clamp API that computes the result with min/max operations [12510]
✅
Out-of-scope changes
Code Change
Explanation
Change extended floating-point fmin and fmax handling (libcudacxx/include/cuda/std/__cmath/min_max.h:70-129,183-226,342-386)
The issue asks for a clamp implementation or callable clamp alternative. It does not state requirements for changing fmin and fmax behavior or implementation.
Add storage-based narrow-format fmin and fmax tests (libcudacxx/test/libcudacxx/std/numerics/c.math/fp_min_max.pass.cpp:202-380)
These tests cover fmin and fmax, not the linked clamp objective.
Suggested reviewers:davebayer, fbusato, jacobfaib
Change: Feature
Merge Risk: Minimal · previous review 4ee52
The previously identified clamp concerns do not remain actionable at this head. No merge-blocking risk is established by the available evidence.
Security architecture risk: Low · up to 4ee52
The inspected changes add value-based arithmetic APIs without introducing a privileged operation or an external service entrypoint. No concrete security issue was identified, but downstream usage and cross-platform behavior are not fully covered.
Retained concerns
No architecture-level concerns identified.
Security review details
Security Blast Radius
inferred — The demonstrated exposure is arithmetic behavior within callers that compile these headers. No tenant, datastore, privileged service, or attacker-reachable downstream sink was established by the supplied relationships or inspected API bodies; maximum application-level exposure remains unresolved because downstream coverage is capped.
Trust Boundaries and Controls
observed — The specifically routed clamp symbols belong to test detection traits, value helpers, assertions, and a test main that uses a locally initialized value rather than command-line input. These ranges do not demonstrate a newly exposed service or an authorization bypass.
- [ ] 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.
Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 11 remain after this review.
While the signature does not need to have const&, the function must not modify the objects passed to it and must be able to accept all values of type (possibly const) Type1 and Type2 regardless of value category (thus, Type1& is not allowed, nor is Type1 unless for Type1 a move is equivalent to a copy(since C++11)).
The types Type1 and Type2 must be such that an object of type T can be implicitly converted to both of them.
Unfortunately, this is related to the Comp functor, not the function signature.
Reviewing files that changed from the base of the PR and between f02c949 and 36fae63.
Files selected for processing: 3
docs/libcudacxx/extended_api/math/clamp.rst
docs/libcudacxx/extended_api/math/min_max.rst
libcudacxx/test/cmath_codegen/clamp.cu
Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 6 remain after this review.
The reason will be displayed to describe this comment to others. Learn more.
Important: The last time we discussed this (#5516, #7325) I strongly argued to name the functions clamp_values etc. to sufficiently distinguish them from their std alternatives taking references. This comment still stands and I would prefer the function to be named clamp_values.
The reason will be displayed to describe this comment to others. Learn more.
I am really not in favor of inventing new names for common operations. We have a lot of standard features with non-standard extensions and we require the user to read the documentation.
This is especially the case here, because the situation where the actual result differs is an incredibly uncommon corner case
sm_75 benchmark for the __nv_bfloat16 min/max change
@bernhardmgruber here are the benchmarks for the three binaries with sm_75 SASS changes.
The T4 run in CI did not produce any data. The T4 stayed at its 585 MHz base clock, so NVBench's throttle check (--throttle-threshold 90, measured against the 1590 MHz boost clock) discarded every sample, and all but one of the 249 workloads timed out without a measurement. The one workload that was measured is not affected by this PR.
So I measured locally. I don't have an sm_75 GPU, so I built the benchmarks as compute_75 PTX only (CMAKE_CUDA_ARCHITECTURES=75-virtual) and ran them on an RTX A6000. The driver JIT-compiles the PTX, so the kernels take the pre-SM80 code path, and CUB selects its sm_75 tuning from the PTX version. Base is the merge-base 5060e9b7d1, test is 6d72810437, with the NVBench flags from ci/bench.yaml, restricted to BF16.
BF16 workload
Base
PR
Change
warp_reduce_min
130.0 µs
37.9 µs
3.4× faster
reduce.min, 2^16
10.2 µs
7.2 µs
1.43× faster
reduce.min, 2^20
14.3 µs
10.2 µs
1.40× faster
reduce.min, 2^24
65.7 µs
60.4 µs
1.09× faster
reduce.min, 2^28
783 µs
790 µs
within noise
transform.pytorchfmin / fmax
within noise at all sizes
The reduce.min rows use 32-bit offsets. 64-bit offsets gave the same results, except 1.07× at 2^24.
The warp benchmark isolates the min operation, so it shows the full effect of replacing the emulation with FMNMX.
The device reduction gains wherever it is not bandwidth-bound. At 2^28 both versions run at 88–89% of peak memory bandwidth.
The PyTorch fmin and fmax transforms are bandwidth-bound at every size. They move by at most one timer tick (1.024 µs), the same as unchanged ops like mul and div.
There are no regressions. This runs the sm_75 code on Ampere hardware, so absolute numbers will differ from a T4. The A6000 has a similar ratio of instruction throughput to memory bandwidth, so I would expect the same trend there.
`cuda::std::min`, `cuda::std::max` and `cuda::std::clamp` must take their
arguments by reference and return a reference to one of them, so we cannot
change them. Add versions to `<cuda/cmath>` that take and return values and
compare floating point values like `cuda::std::fmin` and `cuda::std::fmax`,
so that `cuda::clamp` maps to a single `min` and `max` instruction.
FixesNVIDIA#12510
nvcc with MSVC considers the global `::min` and `::max` overloads of the CUDA
headers when checking calls like `cuda::min(float, double)` in SFINAE contexts,
because `cuda::min` is a single function template in an inline namespace. Add
deleted overloads for mixed argument types, so that such calls select a deleted
function instead of failing deduction.
Also silence `bugprone-branch-clone` in `cuda::clamp`, whose branches are
identical when assertions are disabled.
`cuda::std::isless` handles NaN in constant expressions and supports all
extended floating-point types, so the special cases for them are no longer
needed. Integers are still compared with `<`, because `cuda::std::isless`
would convert them to `double`.
…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.
@bernhardmgruber, your SASS approval is recorded, but GitHub rejects CI rerun requests while it is running. After the current CI run finishes, rerun its failed jobs to apply the approval.
…d `cuda::max`
There should be only one implementation of the minimum and maximum of two
values. `cuda::min` and `cuda::max` take their arguments by value, so the
references that `cuda::minimum` and `cuda::maximum` take are optimized away:
the SASS for `float`, `double`, `__half` and `__nv_bfloat16` does not change.
Integers only differ in the operand order, because `cuda::minimum` compares
`a < b` instead of `b < a`.
- On device, the function maps to a ``min`` and a ``max`` instruction for integer types of up to 64 bits and ``float``, for ``double`` on GPUs with fast FP64 arithmetic, and for ``__half`` and ``__nv_bfloat16`` on ``SM80`` and newer. On GPUs with slow FP64 arithmetic, ``double`` values are compared with integer instructions instead.
⚠️The SASS changed for 3 of 101 CUB and 0 of 75 Thrust benchmark target(s). A benchmark run may be necessary
How to request a benchmark run
Request a benchmark run for this PR:
1. Replace the `benchmarks:` block of ci/bench.yaml with exactly this:
benchmarks:
filters:
cpp:
- '^cub\.bench\.reduce\.min\.base$'
- '^cub\.bench\.reduce\.warp_reduce_min\.base$'
- '^cub\.bench\.transform\.pytorch\.base$'
gpus:
- "h100" # pick the GPUs that this change can affect
2. Commit with `[bench-only]` at the end of the commit summary, so that
the unrelated CI jobs are skipped. Then push.
ci/bench.yaml must match ci/bench.template.yaml before the PR can merge.
Reset it once the measurement is done.
Bfloat16 min and max kernels replace compare branches with FMNMX On sm_75, branched compare, NaN-test, select, and permutation sequences are replaced by FMNMX-based sequences. The signature occurs in one warp-reduce kernel, five device-reduce kernels, and the transform fmin and fmax kernels; instruction counts and control-flow structures change in all three targets.
Showing 40/49166 diff lines, 48949 changes. - ⬇️ Full diff
Bfloat16 min and max kernels replace compare branches with FMNMX On sm_75, branched compare, NaN-test, select, and permutation sequences are replaced by FMNMX-based sequences. The signature occurs in one warp-reduce kernel, five device-reduce kernels, and the transform fmin and fmax kernels; instruction counts and control-flow structures change in all three targets.
Showing 40/166 diff lines, 156 changes. - ⬇️ Full diff
Bfloat16 min and max kernels replace compare branches with FMNMX On sm_75, branched compare, NaN-test, select, and permutation sequences are replaced by FMNMX-based sequences. The signature occurs in one warp-reduce kernel, five device-reduce kernels, and the transform fmin and fmax kernels; instruction counts and control-flow structures change in all three targets.
Showing 40/1865 diff lines, 1768 changes. - ⬇️ Full diff
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
review meetingCandidate for review in live team review meeting.
5 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.
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
davebayer
left a comment
Description
closes #12510
cuda::std::min,cuda::std::maxandcuda::std::clampmust take their arguments by reference and return a reference to one of them, so we cannot change them the way the first version of this PR did. Instead, this addscuda::min,cuda::maxandcuda::clampto<cuda/cmath>, which take and return values. The compiler can then compute the result in registers instead of selecting the address of an argument and loading the result from it, so on devicecuda::clampmaps to a singleminandmaxinstruction. The exception isdoubleon GPUs with slow FP64 arithmetic, wherecuda::std::fminandcuda::std::fmaxcompare with integer instructions since #12231.cuda::std::fminandcuda::std::fmaxsupport since [libcu++] Makefmin/fmaxwork for all nvfp types #12612. All arguments must have the same type.cuda::std::fminandcuda::std::fmax, so a NaN argument is ignored.cuda::clamp(v, lo, hi)iscuda::max(cuda::min(v, hi), lo): a NaN bound is ignored, and a NaNvis clamped tohi, or toloifhiis NaN.cuda::minandcuda::maxcallcuda::minimumandcuda::maximum, so there is a single implementation. Because they take their arguments by value, the references thatcuda::minimumandcuda::maximumtake are optimized away: the SASS forfloat,double,__halfand__nv_bfloat16is the same as when callingcuda::std::fminandcuda::std::fmaxdirectly.Through
cuda::minimumandcuda::maximum, the new functions usecuda::std::fminandcuda::std::fmax, which this also improves:<cuda_fp16.h>and<cuda_bf16.h>are called unqualified and found by ADL, so<cuda/std/__cmath/min_max.h>does not need to include those headers.__nv_bfloat16usesmin.f32andmax.f32instead of the emulation in__hminand__hmax, which takes about 50 instructions with branches. The result is exact, which was verified for all 2^32 input pairs.__hminand__hmaxare only available on device, so the host compares asfloat.Tests and documentation:
cuda/cmath/min_max.pass.cppandcuda/cmath/clamp.pass.cpp, which run with the conversions and operators of the extended floating-point types disabled (__CUDA_NO_HALF_OPERATORS__and friends).cmath-ptxFileCheck suite inlibcudacxx/test/cmath_codegen, which checks thatcuda::min,cuda::maxandcuda::clampcompile tominandmaxinstructions without branches. Fordouble, the instructions are only checked on SM80, SM90 and SM100, because the other targets in CI have slow FP64 arithmetic. It is wired into CI likesimd-ptx.cuda::minandcuda::max, and forcuda::clamp.Checklist