From c6f3878e4eb08ec5abbe7eba46225a6826858b3a Mon Sep 17 00:00:00 2001 From: lennartvoelz Date: Fri, 23 Oct 2026 21:42:33 +0200 Subject: [PATCH 01/13] Add integer path optimization for double prec on slow targets - adds __fp64_to_ordered_int helper that uses a selection form to convert the bit pattern to int64 with the same ordering as fp64 - adds double overloads for isgreater, isless, isgreaterequal, islessequal and islessgreater - uses _CCCL_DISPATCH_SLOW_FP64_TARGET to dispatch between "slow" and "fast" fp64 arithmetic targets --- libcudacxx/include/cuda/std/__cmath/traits.h | 70 ++++++++++++++++++++ 1 file changed, 70 insertions(+) diff --git a/libcudacxx/include/cuda/std/__cmath/traits.h b/libcudacxx/include/cuda/std/__cmath/traits.h index 259124d2cc75..74729c02c1c0 100644 --- a/libcudacxx/include/cuda/std/__cmath/traits.h +++ b/libcudacxx/include/cuda/std/__cmath/traits.h @@ -22,10 +22,13 @@ #endif // no system header #include +#include #include // IWYU pragma: keep #include #include #include +#include +#include #include @@ -33,6 +36,14 @@ _CCCL_BEGIN_NAMESPACE_CUDA_STD +// maps the bits of a double to an int64_t with the same order +// -0 and +0 map to 0, NaNs have to handled by the caller +[[nodiscard]] _CCCL_API inline int64_t __fp64_to_ordered_int(double __x) noexcept +{ + const auto __x_as_int64 = static_cast(::cuda::std::__fp_get_storage(__x)); + return __x_as_int64 < 0 ? ::cuda::std::numeric_limits::min() - __x_as_int64 : __x_as_int64; +} + // isgreater template , int> = 0> @@ -45,6 +56,17 @@ template , int> = 0> return __x > __y; } +[[nodiscard]] _CCCL_API inline bool __device_isgreater(double __x, double __y) noexcept +{ + if (::cuda::std::isnan(__x) || ::cuda::std::isnan(__y)) + { + return false; + } + _CCCL_DISPATCH_SLOW_FP64_TARGET( + return ::cuda::std::__fp64_to_ordered_int(__x) > ::cuda::std::__fp64_to_ordered_int(__y);) + return __x > __y; +} + #if _CCCL_CHECK_BUILTIN(builtin_isgreater) || _CCCL_COMPILER(GCC) # define _CCCL_BUILTIN_ISGREATER(...) __builtin_isgreater(__VA_ARGS__) #endif // _CCCL_CHECK_BUILTIN(builtin_isgreater) @@ -82,6 +104,17 @@ template , int> = 0> return __x >= __y; } +[[nodiscard]] _CCCL_API inline bool __device_isgreaterequal(double __x, double __y) noexcept +{ + if (::cuda::std::isnan(__x) || ::cuda::std::isnan(__y)) + { + return false; + } + _CCCL_DISPATCH_SLOW_FP64_TARGET( + return ::cuda::std::__fp64_to_ordered_int(__x) >= ::cuda::std::__fp64_to_ordered_int(__y);) + return __x >= __y; +} + #if _CCCL_CHECK_BUILTIN(builtin_isgreaterequal) || _CCCL_COMPILER(GCC) # define _CCCL_BUILTIN_ISGREATEREQUAL(...) __builtin_isgreaterequal(__VA_ARGS__) #endif // _CCCL_CHECK_BUILTIN(builtin_isgreaterequal) @@ -119,6 +152,17 @@ template , int> = 0> return __x < __y; } +[[nodiscard]] _CCCL_API inline bool __device_isless(double __x, double __y) noexcept +{ + if (::cuda::std::isnan(__x) || ::cuda::std::isnan(__y)) + { + return false; + } + _CCCL_DISPATCH_SLOW_FP64_TARGET( + return ::cuda::std::__fp64_to_ordered_int(__x) < ::cuda::std::__fp64_to_ordered_int(__y);) + return __x < __y; +} + #if _CCCL_CHECK_BUILTIN(builtin_isless) || _CCCL_COMPILER(GCC) # define _CCCL_BUILTIN_ISLESS(...) __builtin_isless(__VA_ARGS__) #endif // _CCCL_CHECK_BUILTIN(builtin_isless) @@ -156,6 +200,17 @@ template , int> = 0> return __x <= __y; } +[[nodiscard]] _CCCL_API inline bool __device_islessequal(double __x, double __y) noexcept +{ + if (::cuda::std::isnan(__x) || ::cuda::std::isnan(__y)) + { + return false; + } + _CCCL_DISPATCH_SLOW_FP64_TARGET( + return ::cuda::std::__fp64_to_ordered_int(__x) <= ::cuda::std::__fp64_to_ordered_int(__y);) + return __x <= __y; +} + #if _CCCL_CHECK_BUILTIN(builtin_islessequal) || _CCCL_COMPILER(GCC) # define _CCCL_BUILTIN_ISLESSEQUAL(...) __builtin_islessequal(__VA_ARGS__) #endif // _CCCL_CHECK_BUILTIN(builtin_islessequal) @@ -193,6 +248,21 @@ template , int> = 0> return __x < __y || __x > __y; } +[[nodiscard]] _CCCL_API inline bool __device_islessgreater(double __x, double __y) noexcept +{ + if (::cuda::std::isnan(__x) || ::cuda::std::isnan(__y)) + { + return false; + } + // Non-NaN values are unequal iff their bits differ, except for -0 and +0 + // Comparing the bits directly is cheaper than mapping both values to ordered integers + // ((x | y) << 1) == 0 drops the sign bit and is true only if both are +-0 + _CCCL_DISPATCH_SLOW_FP64_TARGET( + const auto __x_bits = ::cuda::std::__fp_get_storage(__x); const auto __y_bits = ::cuda::std::__fp_get_storage(__y); + return __x_bits != __y_bits && ((__x_bits | __y_bits) << 1) != 0;) + return __x < __y || __x > __y; +} + #if _CCCL_CHECK_BUILTIN(builtin_islessgreater) || _CCCL_COMPILER(GCC) # define _CCCL_BUILTIN_ISLESSGREATER(...) __builtin_islessgreater(__VA_ARGS__) #endif // _CCCL_CHECK_BUILTIN(builtin_islessgreater) From 4d87bf04574dda6e2faf902df2d03d04ef57b08b Mon Sep 17 00:00:00 2001 From: lennartvoelz Date: Fri, 23 Oct 2026 21:46:58 +0200 Subject: [PATCH 02/13] Fix typo in comment on __fp64_to_ordered_int --- libcudacxx/include/cuda/std/__cmath/traits.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/libcudacxx/include/cuda/std/__cmath/traits.h b/libcudacxx/include/cuda/std/__cmath/traits.h index 74729c02c1c0..dcf4efe5dd2a 100644 --- a/libcudacxx/include/cuda/std/__cmath/traits.h +++ b/libcudacxx/include/cuda/std/__cmath/traits.h @@ -37,7 +37,7 @@ _CCCL_BEGIN_NAMESPACE_CUDA_STD // maps the bits of a double to an int64_t with the same order -// -0 and +0 map to 0, NaNs have to handled by the caller +// -0 and +0 map to 0, NaNs have to be handled by the caller [[nodiscard]] _CCCL_API inline int64_t __fp64_to_ordered_int(double __x) noexcept { const auto __x_as_int64 = static_cast(::cuda::std::__fp_get_storage(__x)); From 6a471b4b32982a73c2d9d077fa9f00264d032340 Mon Sep 17 00:00:00 2001 From: lennartvoelz Date: Fri, 23 Oct 2026 23:18:37 +0200 Subject: [PATCH 03/13] Fix typo in fmin/fmax comment quite -> quiet in min_max.h --- libcudacxx/include/cuda/std/__cmath/min_max.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/libcudacxx/include/cuda/std/__cmath/min_max.h b/libcudacxx/include/cuda/std/__cmath/min_max.h index e9fc37ea09d0..cc2ea3ea9ae5 100644 --- a/libcudacxx/include/cuda/std/__cmath/min_max.h +++ b/libcudacxx/include/cuda/std/__cmath/min_max.h @@ -47,7 +47,7 @@ _CCCL_REQUIRES(is_floating_point_v<_Tp>) [[nodiscard]] _CCCL_API _Tp __with_builtin_fmax(_Tp __x, _Tp __y) noexcept { # if _CCCL_COMPILER(CLANG) && _CCCL_HOST_ARCH(ARM64) && _CCCL_HOST_COMPILATION() - // Clang lowers fmax to AArch64 fmaxnm, which returns quite NaN for signaling NaNs, + // Clang lowers fmax to AArch64 fmaxnm, which returns quiet NaN for signaling NaNs, // See https://github.com/llvm/llvm-project/issues/176624 if (::cuda::std::isnan(__x)) { @@ -185,7 +185,7 @@ _CCCL_REQUIRES(is_floating_point_v<_Tp>) [[nodiscard]] _CCCL_API _Tp __with_builtin_fmin(_Tp __x, _Tp __y) noexcept { # if _CCCL_COMPILER(CLANG) && _CCCL_HOST_ARCH(ARM64) && _CCCL_HOST_COMPILATION() - // Clang lowers fmin to AArch64 fminnm, which returns quite NaN for signaling NaNs, + // Clang lowers fmin to AArch64 fminnm, which returns quiet NaN for signaling NaNs, // See https://github.com/llvm/llvm-project/issues/176624 if (::cuda::std::isnan(__x)) { From d9758ae4a3472a0c5c3c74ca41fd246cdfcadf84 Mon Sep 17 00:00:00 2001 From: lennartvoelz Date: Sat, 24 Oct 2026 00:40:20 +0200 Subject: [PATCH 04/13] Add integer path optimization for fmin/fmax on slow fp64 targets - adds __fp64_to_total_order_int helper that maps the bit pattern to int64 in IEEE totalOrder (-0 < +0) - adds branch-free __fp64_fmax_fmin_int helper that matches the hardware result bitwise (NaN handling, quiet bit) - uses _CCCL_DISPATCH_SLOW_FP64_TARGET in fmax and fmin to dispatch to the integer path for double --- libcudacxx/include/cuda/std/__cmath/min_max.h | 35 +++++++++++++++++++ 1 file changed, 35 insertions(+) diff --git a/libcudacxx/include/cuda/std/__cmath/min_max.h b/libcudacxx/include/cuda/std/__cmath/min_max.h index cc2ea3ea9ae5..74b6ecc65ca5 100644 --- a/libcudacxx/include/cuda/std/__cmath/min_max.h +++ b/libcudacxx/include/cuda/std/__cmath/min_max.h @@ -24,11 +24,14 @@ #include #include #include +#include #include #include #include #include #include +#include +#include #include @@ -36,6 +39,30 @@ _CCCL_BEGIN_NAMESPACE_CUDA_STD +// maps the bits of a double to an int64_t in IEEE totalOrder: -0 < +0, NaNs are placed beyond +-inf +[[nodiscard]] _CCCL_API constexpr int64_t __fp64_to_total_order_int(double __x) noexcept +{ + const auto __x_as_int = static_cast(::cuda::std::__fp_get_storage(__x)); + return __x_as_int < 0 ? __x_as_int ^ ::cuda::std::numeric_limits::max() : __x_as_int; +} + +// fmax (_IsMax) / fmin for fp64 with integer instructions only +// -0 < +0 is preserved +// single NaN is ignored; two NaNs return __x quieted +template +[[nodiscard]] _CCCL_API constexpr double __fp64_fmax_fmin_int(double __x, double __y) noexcept +{ + constexpr auto __quiet_bit = uint64_t{0x0008'0000'0000'0000}; // top mantissa bit (for quiet NaNs) + const auto __x_bits = ::cuda::std::__fp_get_storage(__x); + const auto __y_bits = ::cuda::std::__fp_get_storage(__y); + const bool __x_nan = ::cuda::std::isnan(__x); + const bool __y_nan = ::cuda::std::isnan(__y); + const bool __x_lt_y = ::cuda::std::__fp64_to_total_order_int(__x) < ::cuda::std::__fp64_to_total_order_int(__y); + const bool __pick_y = __x_nan ? !__y_nan : (!__y_nan && __x_lt_y == _IsMax); + const auto __result = __pick_y ? __y_bits : __x_bits; + return ::cuda::std::__fp_from_storage((__x_nan && __y_nan) ? (__result | __quiet_bit) : __result); +} + /*********************************************************************************************************************** * fmax **********************************************************************************************************************/ @@ -126,6 +153,10 @@ _CCCL_REQUIRES(__is_extended_arithmetic_v<_Tp>) NV_IF_TARGET(NV_PROVIDES_SM_100, (return ::__nv_fp128_fmax(__x, __y);)) } #endif // _CCCL_HAS_FLOAT128() + if constexpr (is_same_v<_Tp, double>) + { + _CCCL_DISPATCH_SLOW_FP64_TARGET(return ::cuda::std::__fp64_fmax_fmin_int(__x, __y);) + } #if _CCCL_USE_BUILTIN_FMAX() if constexpr (is_floating_point_v<_Tp>) { @@ -264,6 +295,10 @@ _CCCL_REQUIRES(__is_extended_arithmetic_v<_Tp>) NV_IF_TARGET(NV_PROVIDES_SM_100, (return ::__nv_fp128_fmin(__x, __y);)) } #endif // _CCCL_HAS_FLOAT128() + if constexpr (is_same_v<_Tp, double>) + { + _CCCL_DISPATCH_SLOW_FP64_TARGET(return ::cuda::std::__fp64_fmax_fmin_int(__x, __y);) + } #if _CCCL_USE_BUILTIN_FMIN() if constexpr (is_floating_point_v<_Tp>) { From c9721c5cd7ec4ac030b9bcdeb0a7883d9dc5c52a Mon Sep 17 00:00:00 2001 From: lennartvoelz Date: Sat, 24 Oct 2026 12:08:46 +0200 Subject: [PATCH 05/13] Add benchmarks for comparison functions and fmin/fmax --- .../benchmarks/bench/comparison/basic.cu | 101 ++++++++++++++++++ .../benchmarks/bench/fmin_fmax/basic.cu | 76 +++++++++++++ 2 files changed, 177 insertions(+) create mode 100644 libcudacxx/benchmarks/bench/comparison/basic.cu create mode 100644 libcudacxx/benchmarks/bench/fmin_fmax/basic.cu diff --git a/libcudacxx/benchmarks/bench/comparison/basic.cu b/libcudacxx/benchmarks/bench/comparison/basic.cu new file mode 100644 index 000000000000..f13dfb287c56 --- /dev/null +++ b/libcudacxx/benchmarks/bench/comparison/basic.cu @@ -0,0 +1,101 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#include +#include + +#include +#include + +#include + +struct double_pair +{ + double x{}; + double y{}; +}; + +struct isless_op +{ + __device__ __forceinline__ bool operator()(double x, double y) const + { + return cuda::std::isless(x, y); + } +}; + +struct islessequal_op +{ + __device__ __forceinline__ bool operator()(double x, double y) const + { + return cuda::std::islessequal(x, y); + } +}; + +struct isgreater_op +{ + __device__ __forceinline__ bool operator()(double x, double y) const + { + return cuda::std::isgreater(x, y); + } +}; + +struct isgreaterequal_op +{ + __device__ __forceinline__ bool operator()(double x, double y) const + { + return cuda::std::isgreaterequal(x, y); + } +}; + +struct islessgreater_op +{ + __device__ __forceinline__ bool operator()(double x, double y) const + { + return cuda::std::islessgreater(x, y); + } +}; + +template +struct benchmark_op_t +{ + __device__ __forceinline__ double_pair operator()(double_pair value) const + { + // both operands change every iteration, so the compiler cannot hoist work on a constant operand + // storage is used to prevent code elimination by the compiler (same approach as in isnan) + // __fp_get_storage and __fp_from_storage are no op + const auto storage = cuda::std::__fp_get_storage(value.x); + return {value.y, cuda::std::__fp_from_storage(storage + CompareOp{}(value.x, value.y))}; + } +}; + +template +void comparison_bench(nvbench::state& state, nvbench::type_list) +{ + constexpr int block_size = 256; + constexpr int unroll_factor = 1024; + const auto& kernel = benchmark_kernel, double_pair, false>; + const int num_SMs = state.get_device().value().get_number_of_sms(); // NOLINT(bugprone-unchecked-optional-access) + int max_blocks_per_SM = 0; + NVBENCH_CUDA_CALL_NOEXCEPT(cudaOccupancyMaxActiveBlocksPerMultiprocessor(&max_blocks_per_SM, kernel, block_size, 0)); + const int grid_size = max_blocks_per_SM * num_SMs; + state.exec(nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch, [&](nvbench::launch&) { + kernel<<>>(benchmark_op_t{}); + }); +} + +NVBENCH_DECLARE_TYPE_STRINGS(isless_op, "isless", ""); +NVBENCH_DECLARE_TYPE_STRINGS(islessequal_op, "islessequal", ""); +NVBENCH_DECLARE_TYPE_STRINGS(isgreater_op, "isgreater", ""); +NVBENCH_DECLARE_TYPE_STRINGS(isgreaterequal_op, "isgreaterequal", ""); +NVBENCH_DECLARE_TYPE_STRINGS(islessgreater_op, "islessgreater", ""); + +using ops = nvbench::type_list; + +NVBENCH_BENCH_TYPES(comparison_bench, NVBENCH_TYPE_AXES(ops)).set_name("base").set_type_axes_names({"op"}); diff --git a/libcudacxx/benchmarks/bench/fmin_fmax/basic.cu b/libcudacxx/benchmarks/bench/fmin_fmax/basic.cu new file mode 100644 index 000000000000..646572bbdf25 --- /dev/null +++ b/libcudacxx/benchmarks/bench/fmin_fmax/basic.cu @@ -0,0 +1,76 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#include +#include + +#include +#include + +#include + +struct double_pair +{ + double x{}; + double y{}; +}; + +struct fmax_op +{ + __device__ __forceinline__ double operator()(double x, double y) const + { + return cuda::std::fmax(x, y); + } +}; + +struct fmin_op +{ + __device__ __forceinline__ double operator()(double x, double y) const + { + return cuda::std::fmin(x, y); + } +}; + +template +struct benchmark_op_t +{ + __device__ __forceinline__ double_pair operator()(double_pair value) const + { + // both operands change every iteration, so the compiler cannot hoist work on a constant operand + // storage is used to prevent code elimination by the compiler (adapted from isnan benchmark) + // the result is mixed with the bits of x, otherwise the chain settles on a single value (x == y) + // __fp_get_storage and __fp_from_storage are no op + const auto storage = cuda::std::__fp_get_storage(value.x); + const auto result = cuda::std::__fp_get_storage(MinMaxOp{}(value.x, value.y)); + return {value.y, cuda::std::__fp_from_storage(result ^ (storage >> 7))}; + } +}; + +template +void min_max_bench(nvbench::state& state, nvbench::type_list) +{ + constexpr int block_size = 256; + constexpr int unroll_factor = 1024; + const auto& kernel = benchmark_kernel, double_pair, false>; + const int num_SMs = state.get_device().value().get_number_of_sms(); // NOLINT(bugprone-unchecked-optional-access) + int max_blocks_per_SM = 0; + NVBENCH_CUDA_CALL_NOEXCEPT(cudaOccupancyMaxActiveBlocksPerMultiprocessor(&max_blocks_per_SM, kernel, block_size, 0)); + const int grid_size = max_blocks_per_SM * num_SMs; + state.exec(nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch, [&](nvbench::launch&) { + kernel<<>>(benchmark_op_t{}); + }); +} + +NVBENCH_DECLARE_TYPE_STRINGS(fmax_op, "fmax", ""); +NVBENCH_DECLARE_TYPE_STRINGS(fmin_op, "fmin", ""); + +using ops = nvbench::type_list; + +NVBENCH_BENCH_TYPES(min_max_bench, NVBENCH_TYPE_AXES(ops)).set_name("base").set_type_axes_names({"op"}); From 64fa02e423294179b75e71285fc5783245ec506b Mon Sep 17 00:00:00 2001 From: lennartvoelz Date: Sat, 24 Oct 2026 13:28:28 +0200 Subject: [PATCH 06/13] Make the comparison functions branch-free -> another 40% performance gain --- libcudacxx/include/cuda/std/__cmath/traits.h | 24 +++++++++++++------- 1 file changed, 16 insertions(+), 8 deletions(-) diff --git a/libcudacxx/include/cuda/std/__cmath/traits.h b/libcudacxx/include/cuda/std/__cmath/traits.h index dcf4efe5dd2a..cd648073d8ca 100644 --- a/libcudacxx/include/cuda/std/__cmath/traits.h +++ b/libcudacxx/include/cuda/std/__cmath/traits.h @@ -58,12 +58,14 @@ template , int> = 0> [[nodiscard]] _CCCL_API inline bool __device_isgreater(double __x, double __y) noexcept { + // compute comparison before NaN check, so compiler emits no branch around it + _CCCL_DISPATCH_SLOW_FP64_TARGET( + const bool __is_greater = ::cuda::std::__fp64_to_ordered_int(__x) > ::cuda::std::__fp64_to_ordered_int(__y); + return !(::cuda::std::isnan(__x) || ::cuda::std::isnan(__y)) && __is_greater;) if (::cuda::std::isnan(__x) || ::cuda::std::isnan(__y)) { return false; } - _CCCL_DISPATCH_SLOW_FP64_TARGET( - return ::cuda::std::__fp64_to_ordered_int(__x) > ::cuda::std::__fp64_to_ordered_int(__y);) return __x > __y; } @@ -106,12 +108,14 @@ template , int> = 0> [[nodiscard]] _CCCL_API inline bool __device_isgreaterequal(double __x, double __y) noexcept { + // compute comparison before NaN check, so compiler emits no branch around it + _CCCL_DISPATCH_SLOW_FP64_TARGET( + const bool __is_greater_equal = ::cuda::std::__fp64_to_ordered_int(__x) >= ::cuda::std::__fp64_to_ordered_int(__y); + return !(::cuda::std::isnan(__x) || ::cuda::std::isnan(__y)) && __is_greater_equal;) if (::cuda::std::isnan(__x) || ::cuda::std::isnan(__y)) { return false; } - _CCCL_DISPATCH_SLOW_FP64_TARGET( - return ::cuda::std::__fp64_to_ordered_int(__x) >= ::cuda::std::__fp64_to_ordered_int(__y);) return __x >= __y; } @@ -154,12 +158,14 @@ template , int> = 0> [[nodiscard]] _CCCL_API inline bool __device_isless(double __x, double __y) noexcept { + // compute comparison before NaN check, so compiler emits no branch around it + _CCCL_DISPATCH_SLOW_FP64_TARGET( + const bool __is_less = ::cuda::std::__fp64_to_ordered_int(__x) < ::cuda::std::__fp64_to_ordered_int(__y); + return !(::cuda::std::isnan(__x) || ::cuda::std::isnan(__y)) && __is_less;) if (::cuda::std::isnan(__x) || ::cuda::std::isnan(__y)) { return false; } - _CCCL_DISPATCH_SLOW_FP64_TARGET( - return ::cuda::std::__fp64_to_ordered_int(__x) < ::cuda::std::__fp64_to_ordered_int(__y);) return __x < __y; } @@ -202,12 +208,14 @@ template , int> = 0> [[nodiscard]] _CCCL_API inline bool __device_islessequal(double __x, double __y) noexcept { + // compute comparison before NaN check, so compiler emits no branch around it + _CCCL_DISPATCH_SLOW_FP64_TARGET( + const bool __is_less_equal = ::cuda::std::__fp64_to_ordered_int(__x) <= ::cuda::std::__fp64_to_ordered_int(__y); + return !(::cuda::std::isnan(__x) || ::cuda::std::isnan(__y)) && __is_less_equal;) if (::cuda::std::isnan(__x) || ::cuda::std::isnan(__y)) { return false; } - _CCCL_DISPATCH_SLOW_FP64_TARGET( - return ::cuda::std::__fp64_to_ordered_int(__x) <= ::cuda::std::__fp64_to_ordered_int(__y);) return __x <= __y; } From 1dda4b5b5ca48aca6591203aae2b9c9a0fd54f60 Mon Sep 17 00:00:00 2001 From: lennartvoelz Date: Mon, 26 Oct 2026 07:18:34 +0100 Subject: [PATCH 07/13] Add edge case tests for double on slow FP64 targets - tests for fmin/fmax and the comparison functions isless, islessequal, isgreater, isgreaterequal, islessgreater --- .../std/numerics/c.math/fp_min_max.pass.cpp | 83 +++++++++++++++++++ .../std/numerics/c.math/fp_traits.pass.cpp | 70 ++++++++++++++++ 2 files changed, 153 insertions(+) diff --git a/libcudacxx/test/libcudacxx/std/numerics/c.math/fp_min_max.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/c.math/fp_min_max.pass.cpp index 84d10dfc43cd..b7f97687c543 100644 --- a/libcudacxx/test/libcudacxx/std/numerics/c.math/fp_min_max.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/numerics/c.math/fp_min_max.pass.cpp @@ -10,10 +10,15 @@ // #include +#include #include #include +#include +#include #include +#include + #include "comparison.h" #include "test_macros.h" @@ -106,6 +111,82 @@ TEST_FUNC constexpr void test_fmin(T val) } } +TEST_FUNC void test_fp64_min_max(double x, double y, double max, double min) +{ + assert(eq(cuda::std::fmax(x, y), max)); + assert(eq(cuda::std::fmin(x, y), min)); +} + +// double has an integer implementation on slow fp64 targets +// check the cases where the integer mapping can go wrong +TEST_FUNC void test_fp64_edge_cases(float val) +{ + // values are derived from val at runtime + double one = val; + double two = one + one; + double zero = one - one; + double nzero = -zero; + double inf = cuda::std::numeric_limits::infinity() * one; + double max = cuda::std::numeric_limits::max() * one; + double dmin = cuda::std::numeric_limits::denorm_min() * one; + double nan = cuda::std::numeric_limits::quiet_NaN(); + double nnan = cuda::std::copysign(nan, -one); + // nvcc quiets signaling NaN constants in device code, so the bits are built at runtime + double snan = + cuda::std::bit_cast(cuda::std::uint64_t{0x7ff4'0000'0000'0000} * static_cast(val)); + + // x y fmax fmin + // signed zeros compare equal + // only the value is checked + test_fp64_min_max(zero, nzero, zero, zero); + test_fp64_min_max(nzero, zero, zero, zero); + assert(cuda::std::signbit(cuda::std::fmax(nzero, nzero))); + assert(cuda::std::signbit(cuda::std::fmin(nzero, nzero))); + + // negative values are ordered in reverse to their bit patterns + test_fp64_min_max(-one, -two, -one, -two); + test_fp64_min_max(-two, -one, -one, -two); + test_fp64_min_max(-one, one, one, -one); + test_fp64_min_max(-one, -one, -one, -one); + + // infinities and limits + test_fp64_min_max(-inf, -max, -max, -inf); + test_fp64_min_max(max, inf, inf, max); + + // denormals + test_fp64_min_max(dmin, zero, dmin, zero); + test_fp64_min_max(-dmin, nzero, nzero, -dmin); + test_fp64_min_max(-dmin, dmin, dmin, -dmin); + + // negative NaNs map below -inf, but the other value must still be returned + test_fp64_min_max(nnan, -one, -one, -one); + test_fp64_min_max(-one, nnan, -one, -one); + test_fp64_min_max(nnan, -inf, -inf, -inf); + test_fp64_min_max(snan, one, one, one); + test_fp64_min_max(one, snan, one, one); + + // two NaNs return a NaN, the payload and whether a signaling NaN is quieted are unspecified + assert(cuda::std::isnan(cuda::std::fmax(nan, nan))); + assert(cuda::std::isnan(cuda::std::fmin(nan, nan))); + assert(cuda::std::isnan(cuda::std::fmax(nnan, nan))); + assert(cuda::std::isnan(cuda::std::fmin(nnan, nan))); + assert(cuda::std::isnan(cuda::std::fmax(snan, snan))); + assert(cuda::std::isnan(cuda::std::fmin(snan, snan))); + assert(cuda::std::isnan(cuda::std::fmax(nan, snan))); + assert(cuda::std::isnan(cuda::std::fmin(snan, nan))); + + // mixed arguments are promoted to double + assert(eq(cuda::std::fmax(-val, -0.5 * one), -0.5)); + assert(eq(cuda::std::fmin(-val, -0.5 * one), -1.0)); + + // on device -0 < +0 for fmax and fmin + NV_IF_TARGET(NV_IS_DEVICE, + (assert(!cuda::std::signbit(cuda::std::fmax(zero, nzero))); + assert(!cuda::std::signbit(cuda::std::fmax(nzero, zero))); + assert(cuda::std::signbit(cuda::std::fmin(zero, nzero))); + assert(cuda::std::signbit(cuda::std::fmin(nzero, zero)));)) +} + template TEST_FUNC constexpr bool test(T val) { @@ -148,6 +229,8 @@ TEST_FUNC bool test(float val) test<__uint128_t>(static_cast<__uint128_t>(static_cast(val))); #endif // _CCCL_HAS_INT128() + test_fp64_edge_cases(val); + return true; } diff --git a/libcudacxx/test/libcudacxx/std/numerics/c.math/fp_traits.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/c.math/fp_traits.pass.cpp index 79fc48074009..b98a56f16884 100644 --- a/libcudacxx/test/libcudacxx/std/numerics/c.math/fp_traits.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/numerics/c.math/fp_traits.pass.cpp @@ -230,6 +230,75 @@ TEST_FUNC void test_isunordered(float val) #endif // _CCCL_HAS_FLOAT128() } +TEST_FUNC void test_fp64_comparison( + double x, double y, bool less, bool less_equal, bool greater, bool greater_equal, bool less_greater) +{ + assert(cuda::std::isless(x, y) == less); + assert(cuda::std::islessequal(x, y) == less_equal); + assert(cuda::std::isgreater(x, y) == greater); + assert(cuda::std::isgreaterequal(x, y) == greater_equal); + assert(cuda::std::islessgreater(x, y) == less_greater); +} + +// double has an integer implementation on slow fp64 targets +// check the cases where the integer mapping can go wrong +TEST_FUNC void test_fp64_edge_cases(float val) +{ + // values are derived from val at runtime + double one = val; + double two = one + one; + double zero = one - one; + double nzero = -zero; + double inf = cuda::std::numeric_limits::infinity() * one; + double max = cuda::std::numeric_limits::max() * one; + double dmin = cuda::std::numeric_limits::denorm_min() * one; + double nan = cuda::std::numeric_limits::quiet_NaN(); + double nnan = cuda::std::copysign(nan, -one); + + // x y < <= > >= <> + // signed zeros compare equal + test_fp64_comparison(zero, nzero, false, true, false, true, false); + test_fp64_comparison(nzero, zero, false, true, false, true, false); + + // negative values are ordered in reverse to their bit patterns + test_fp64_comparison(-one, -two, false, false, true, true, true); + test_fp64_comparison(-two, -one, true, true, false, false, true); + test_fp64_comparison(-one, one, true, true, false, false, true); + test_fp64_comparison(one, -one, false, false, true, true, true); + + // equal values + test_fp64_comparison(one, one, false, true, false, true, false); + test_fp64_comparison(-one, -one, false, true, false, true, false); + + // infinities and limits + test_fp64_comparison(-inf, inf, true, true, false, false, true); + test_fp64_comparison(inf, inf, false, true, false, true, false); + test_fp64_comparison(-inf, -max, true, true, false, false, true); + test_fp64_comparison(max, inf, true, true, false, false, true); + + // denormals + test_fp64_comparison(dmin, zero, false, false, true, true, true); + test_fp64_comparison(-dmin, zero, true, true, false, false, true); + test_fp64_comparison(-dmin, nzero, true, true, false, false, true); + test_fp64_comparison(-dmin, dmin, true, true, false, false, true); + + // NaNs are unordered, including negative NaNs whose bit patterns would map below -inf + test_fp64_comparison(nan, one, false, false, false, false, false); + test_fp64_comparison(one, nan, false, false, false, false, false); + test_fp64_comparison(nan, nan, false, false, false, false, false); + test_fp64_comparison(nan, inf, false, false, false, false, false); + test_fp64_comparison(nnan, -one, false, false, false, false, false); + test_fp64_comparison(-inf, nnan, false, false, false, false, false); + + // mixed arguments are promoted to double + int int_one = static_cast(val); + assert(cuda::std::isless(-int_one, -0.5 * one) == true); + assert(cuda::std::isgreater(int_one - int_one, nzero) == false); + assert(cuda::std::islessequal(int_one - int_one, nzero) == true); + assert(cuda::std::islessgreater(-2 * int_one, -one) == true); + assert(cuda::std::isless(-val, -0.5 * one) == true); +} + TEST_FUNC void test(float val) { test_isgreater(val); @@ -238,6 +307,7 @@ TEST_FUNC void test(float val) test_islessequal(val); test_islessgreater(val); test_isunordered(val); + test_fp64_edge_cases(val); } __global__ void test_global_kernel(float* val) From 5cb1723b76957cde124c891b0ec3fe63cd44e591 Mon Sep 17 00:00:00 2001 From: lennartvoelz Date: Mon, 26 Oct 2026 07:34:52 +0100 Subject: [PATCH 08/13] =?UTF-8?q?=20Under=20NVRTC=20libcudac++=20does=20no?= =?UTF-8?q?t=20use=20fmax/fmin=20builtins,=20so=20on=20fast=20targets=20th?= =?UTF-8?q?e=20generic=20fallback=20(x=20<=20y=20=3F=20y=20:=20x)=20return?= =?UTF-8?q?s=20either=20zero=20for=20=C2=B10,=20while=20the=20integer=20pa?= =?UTF-8?q?th=20on=20slow=20targets=20orders=20-0=20<=20+0?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- .../test/libcudacxx/std/numerics/c.math/fp_min_max.pass.cpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/libcudacxx/test/libcudacxx/std/numerics/c.math/fp_min_max.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/c.math/fp_min_max.pass.cpp index b7f97687c543..d9ab7c9628cb 100644 --- a/libcudacxx/test/libcudacxx/std/numerics/c.math/fp_min_max.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/numerics/c.math/fp_min_max.pass.cpp @@ -180,11 +180,14 @@ TEST_FUNC void test_fp64_edge_cases(float val) assert(eq(cuda::std::fmin(-val, -0.5 * one), -1.0)); // on device -0 < +0 for fmax and fmin + // libcu++ does not use the fmax/fmin builtins under NVRTC, the fallback on fast fp64 targets may return either zero +#if !TEST_COMPILER(NVRTC) NV_IF_TARGET(NV_IS_DEVICE, (assert(!cuda::std::signbit(cuda::std::fmax(zero, nzero))); assert(!cuda::std::signbit(cuda::std::fmax(nzero, zero))); assert(cuda::std::signbit(cuda::std::fmin(zero, nzero))); assert(cuda::std::signbit(cuda::std::fmin(nzero, zero)));)) +#endif // !TEST_COMPILER(NVRTC) } template From 75318ffb96787863a785f25ca0182f5acbab1160 Mon Sep 17 00:00:00 2001 From: lennartvoelz Date: Mon, 26 Oct 2026 13:26:56 +0100 Subject: [PATCH 09/13] Use non-extended internal limits header --- libcudacxx/include/cuda/std/__cmath/min_max.h | 2 +- libcudacxx/include/cuda/std/__cmath/traits.h | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/libcudacxx/include/cuda/std/__cmath/min_max.h b/libcudacxx/include/cuda/std/__cmath/min_max.h index 74b6ecc65ca5..64f24c107060 100644 --- a/libcudacxx/include/cuda/std/__cmath/min_max.h +++ b/libcudacxx/include/cuda/std/__cmath/min_max.h @@ -25,13 +25,13 @@ #include #include #include +#include #include #include #include #include #include #include -#include #include diff --git a/libcudacxx/include/cuda/std/__cmath/traits.h b/libcudacxx/include/cuda/std/__cmath/traits.h index cd648073d8ca..e8cfb4063b43 100644 --- a/libcudacxx/include/cuda/std/__cmath/traits.h +++ b/libcudacxx/include/cuda/std/__cmath/traits.h @@ -24,11 +24,11 @@ #include #include #include // IWYU pragma: keep +#include #include #include #include #include -#include #include From 9fa59493cbfc730725fb264e2c67e32899479709 Mon Sep 17 00:00:00 2001 From: lennartvoelz Date: Mon, 26 Oct 2026 18:25:51 +0100 Subject: [PATCH 10/13] Wrap multi-line statements in FP64 slow target macroi body in braces --- libcudacxx/include/cuda/std/__cmath/traits.h | 28 ++++++++++++-------- 1 file changed, 17 insertions(+), 11 deletions(-) diff --git a/libcudacxx/include/cuda/std/__cmath/traits.h b/libcudacxx/include/cuda/std/__cmath/traits.h index e8cfb4063b43..02b54ac59051 100644 --- a/libcudacxx/include/cuda/std/__cmath/traits.h +++ b/libcudacxx/include/cuda/std/__cmath/traits.h @@ -59,9 +59,10 @@ template , int> = 0> [[nodiscard]] _CCCL_API inline bool __device_isgreater(double __x, double __y) noexcept { // compute comparison before NaN check, so compiler emits no branch around it - _CCCL_DISPATCH_SLOW_FP64_TARGET( + _CCCL_DISPATCH_SLOW_FP64_TARGET({ const bool __is_greater = ::cuda::std::__fp64_to_ordered_int(__x) > ::cuda::std::__fp64_to_ordered_int(__y); - return !(::cuda::std::isnan(__x) || ::cuda::std::isnan(__y)) && __is_greater;) + return !(::cuda::std::isnan(__x) || ::cuda::std::isnan(__y)) && __is_greater; + }) if (::cuda::std::isnan(__x) || ::cuda::std::isnan(__y)) { return false; @@ -109,9 +110,10 @@ template , int> = 0> [[nodiscard]] _CCCL_API inline bool __device_isgreaterequal(double __x, double __y) noexcept { // compute comparison before NaN check, so compiler emits no branch around it - _CCCL_DISPATCH_SLOW_FP64_TARGET( + _CCCL_DISPATCH_SLOW_FP64_TARGET({ const bool __is_greater_equal = ::cuda::std::__fp64_to_ordered_int(__x) >= ::cuda::std::__fp64_to_ordered_int(__y); - return !(::cuda::std::isnan(__x) || ::cuda::std::isnan(__y)) && __is_greater_equal;) + return !(::cuda::std::isnan(__x) || ::cuda::std::isnan(__y)) && __is_greater_equal; + }) if (::cuda::std::isnan(__x) || ::cuda::std::isnan(__y)) { return false; @@ -159,9 +161,10 @@ template , int> = 0> [[nodiscard]] _CCCL_API inline bool __device_isless(double __x, double __y) noexcept { // compute comparison before NaN check, so compiler emits no branch around it - _CCCL_DISPATCH_SLOW_FP64_TARGET( + _CCCL_DISPATCH_SLOW_FP64_TARGET({ const bool __is_less = ::cuda::std::__fp64_to_ordered_int(__x) < ::cuda::std::__fp64_to_ordered_int(__y); - return !(::cuda::std::isnan(__x) || ::cuda::std::isnan(__y)) && __is_less;) + return !(::cuda::std::isnan(__x) || ::cuda::std::isnan(__y)) && __is_less; + }) if (::cuda::std::isnan(__x) || ::cuda::std::isnan(__y)) { return false; @@ -209,9 +212,10 @@ template , int> = 0> [[nodiscard]] _CCCL_API inline bool __device_islessequal(double __x, double __y) noexcept { // compute comparison before NaN check, so compiler emits no branch around it - _CCCL_DISPATCH_SLOW_FP64_TARGET( + _CCCL_DISPATCH_SLOW_FP64_TARGET({ const bool __is_less_equal = ::cuda::std::__fp64_to_ordered_int(__x) <= ::cuda::std::__fp64_to_ordered_int(__y); - return !(::cuda::std::isnan(__x) || ::cuda::std::isnan(__y)) && __is_less_equal;) + return !(::cuda::std::isnan(__x) || ::cuda::std::isnan(__y)) && __is_less_equal; + }) if (::cuda::std::isnan(__x) || ::cuda::std::isnan(__y)) { return false; @@ -265,9 +269,11 @@ template , int> = 0> // Non-NaN values are unequal iff their bits differ, except for -0 and +0 // Comparing the bits directly is cheaper than mapping both values to ordered integers // ((x | y) << 1) == 0 drops the sign bit and is true only if both are +-0 - _CCCL_DISPATCH_SLOW_FP64_TARGET( - const auto __x_bits = ::cuda::std::__fp_get_storage(__x); const auto __y_bits = ::cuda::std::__fp_get_storage(__y); - return __x_bits != __y_bits && ((__x_bits | __y_bits) << 1) != 0;) + _CCCL_DISPATCH_SLOW_FP64_TARGET({ + const auto __x_bits = ::cuda::std::__fp_get_storage(__x); + const auto __y_bits = ::cuda::std::__fp_get_storage(__y); + return __x_bits != __y_bits && ((__x_bits | __y_bits) << 1) != 0; + }) return __x < __y || __x > __y; } From 344acb716c853eba7bb32dec53bb7a1aae6a602d Mon Sep 17 00:00:00 2001 From: lennartvoelz Date: Tue, 27 Oct 2026 05:52:36 +0100 Subject: [PATCH 11/13] Return fixed quiet NaN for two NaN inputs in FP64 fmin/fmax integer path --- libcudacxx/include/cuda/std/__cmath/min_max.h | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/libcudacxx/include/cuda/std/__cmath/min_max.h b/libcudacxx/include/cuda/std/__cmath/min_max.h index 64f24c107060..54468c8ce86c 100644 --- a/libcudacxx/include/cuda/std/__cmath/min_max.h +++ b/libcudacxx/include/cuda/std/__cmath/min_max.h @@ -48,11 +48,11 @@ _CCCL_BEGIN_NAMESPACE_CUDA_STD // fmax (_IsMax) / fmin for fp64 with integer instructions only // -0 < +0 is preserved -// single NaN is ignored; two NaNs return __x quieted +// single NaN is ignored; two NaNs return a fixed quiet NaN template [[nodiscard]] _CCCL_API constexpr double __fp64_fmax_fmin_int(double __x, double __y) noexcept { - constexpr auto __quiet_bit = uint64_t{0x0008'0000'0000'0000}; // top mantissa bit (for quiet NaNs) + constexpr auto __quiet_nan = uint64_t{0x7FF8'0000'0000'0000}; // fixed quiet NaN const auto __x_bits = ::cuda::std::__fp_get_storage(__x); const auto __y_bits = ::cuda::std::__fp_get_storage(__y); const bool __x_nan = ::cuda::std::isnan(__x); @@ -60,7 +60,7 @@ template const bool __x_lt_y = ::cuda::std::__fp64_to_total_order_int(__x) < ::cuda::std::__fp64_to_total_order_int(__y); const bool __pick_y = __x_nan ? !__y_nan : (!__y_nan && __x_lt_y == _IsMax); const auto __result = __pick_y ? __y_bits : __x_bits; - return ::cuda::std::__fp_from_storage((__x_nan && __y_nan) ? (__result | __quiet_bit) : __result); + return ::cuda::std::__fp_from_storage((__x_nan && __y_nan) ? __quiet_nan : __result); } /*********************************************************************************************************************** From 7cf03277b627aae7af8831848239955b878af365 Mon Sep 17 00:00:00 2001 From: lennartvoelz Date: Tue, 27 Oct 2026 07:21:58 +0100 Subject: [PATCH 12/13] Remove magic constant; use numeric limits API instead --- libcudacxx/include/cuda/std/__cmath/min_max.h | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/libcudacxx/include/cuda/std/__cmath/min_max.h b/libcudacxx/include/cuda/std/__cmath/min_max.h index 54468c8ce86c..90e427300c39 100644 --- a/libcudacxx/include/cuda/std/__cmath/min_max.h +++ b/libcudacxx/include/cuda/std/__cmath/min_max.h @@ -52,15 +52,15 @@ _CCCL_BEGIN_NAMESPACE_CUDA_STD template [[nodiscard]] _CCCL_API constexpr double __fp64_fmax_fmin_int(double __x, double __y) noexcept { - constexpr auto __quiet_nan = uint64_t{0x7FF8'0000'0000'0000}; // fixed quiet NaN - const auto __x_bits = ::cuda::std::__fp_get_storage(__x); - const auto __y_bits = ::cuda::std::__fp_get_storage(__y); - const bool __x_nan = ::cuda::std::isnan(__x); - const bool __y_nan = ::cuda::std::isnan(__y); + const auto __x_bits = ::cuda::std::__fp_get_storage(__x); + const auto __y_bits = ::cuda::std::__fp_get_storage(__y); + const bool __x_nan = ::cuda::std::isnan(__x); + const bool __y_nan = ::cuda::std::isnan(__y); const bool __x_lt_y = ::cuda::std::__fp64_to_total_order_int(__x) < ::cuda::std::__fp64_to_total_order_int(__y); const bool __pick_y = __x_nan ? !__y_nan : (!__y_nan && __x_lt_y == _IsMax); const auto __result = __pick_y ? __y_bits : __x_bits; - return ::cuda::std::__fp_from_storage((__x_nan && __y_nan) ? __quiet_nan : __result); + return ::cuda::std::__fp_from_storage( + (__x_nan && __y_nan) ? ::cuda::std::__fp_get_storage(::cuda::std::numeric_limits::quiet_NaN()) : __result); } /*********************************************************************************************************************** From c3f837cbe9cb7eea006cf72e26bfa5a2310d87ba Mon Sep 17 00:00:00 2001 From: lennartvoelz Date: Thu, 29 Oct 2026 00:07:50 +0100 Subject: [PATCH 13/13] Split long quiet NaN expression in return statement in fmin/fmax ineger path --- libcudacxx/include/cuda/std/__cmath/min_max.h | 18 +++++++++--------- 1 file changed, 9 insertions(+), 9 deletions(-) diff --git a/libcudacxx/include/cuda/std/__cmath/min_max.h b/libcudacxx/include/cuda/std/__cmath/min_max.h index 90e427300c39..a0efccbf6310 100644 --- a/libcudacxx/include/cuda/std/__cmath/min_max.h +++ b/libcudacxx/include/cuda/std/__cmath/min_max.h @@ -52,15 +52,15 @@ _CCCL_BEGIN_NAMESPACE_CUDA_STD template [[nodiscard]] _CCCL_API constexpr double __fp64_fmax_fmin_int(double __x, double __y) noexcept { - const auto __x_bits = ::cuda::std::__fp_get_storage(__x); - const auto __y_bits = ::cuda::std::__fp_get_storage(__y); - const bool __x_nan = ::cuda::std::isnan(__x); - const bool __y_nan = ::cuda::std::isnan(__y); - const bool __x_lt_y = ::cuda::std::__fp64_to_total_order_int(__x) < ::cuda::std::__fp64_to_total_order_int(__y); - const bool __pick_y = __x_nan ? !__y_nan : (!__y_nan && __x_lt_y == _IsMax); - const auto __result = __pick_y ? __y_bits : __x_bits; - return ::cuda::std::__fp_from_storage( - (__x_nan && __y_nan) ? ::cuda::std::__fp_get_storage(::cuda::std::numeric_limits::quiet_NaN()) : __result); + const auto __x_bits = ::cuda::std::__fp_get_storage(__x); + const auto __y_bits = ::cuda::std::__fp_get_storage(__y); + const bool __x_nan = ::cuda::std::isnan(__x); + const bool __y_nan = ::cuda::std::isnan(__y); + const bool __x_lt_y = ::cuda::std::__fp64_to_total_order_int(__x) < ::cuda::std::__fp64_to_total_order_int(__y); + const bool __pick_y = __x_nan ? !__y_nan : (!__y_nan && __x_lt_y == _IsMax); + const auto __result = __pick_y ? __y_bits : __x_bits; + const auto __quiet_bits = ::cuda::std::__fp_get_storage(::cuda::std::numeric_limits::quiet_NaN()); + return ::cuda::std::__fp_from_storage((__x_nan && __y_nan) ? __quiet_bits : __result); } /***********************************************************************************************************************