From 8b377e6568523e8615420a699d1d4458691f66a1 Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Fri, 8 Nov 2024 14:20:22 +0100 Subject: [PATCH] Move definition of logarithms into their own file and implement them on our own Fiyes #2743 --- .../include/cuda/std/__cmath/logarithms.h | 506 ++++++++++++++++++ libcudacxx/include/cuda/std/__cmath/nvbf16.h | 6 +- libcudacxx/include/cuda/std/__cmath/nvfp16.h | 28 +- .../cuda/std/detail/libcxx/include/cmath | 24 +- .../std/numerics/c.math/logarithms.pass.cpp | 205 +++++++ 5 files changed, 714 insertions(+), 55 deletions(-) create mode 100644 libcudacxx/include/cuda/std/__cmath/logarithms.h create mode 100644 libcudacxx/test/libcudacxx/std/numerics/c.math/logarithms.pass.cpp diff --git a/libcudacxx/include/cuda/std/__cmath/logarithms.h b/libcudacxx/include/cuda/std/__cmath/logarithms.h new file mode 100644 index 00000000000..913c92f7c57 --- /dev/null +++ b/libcudacxx/include/cuda/std/__cmath/logarithms.h @@ -0,0 +1,506 @@ +// -*- C++ -*- +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, 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) 2024 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _LIBCUDACXX___CMATH_LOGARITHMS_H +#define _LIBCUDACXX___CMATH_LOGARITHMS_H + +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +#include +#include +#include +#include + +#include + +// MSVC does not know the builtins. +// clang-cuda fails with an unresolved symbol +#if defined(_CCCL_COMPILER_MSVC) || defined(_CCCL_CUDA_COMPILER_CLANG) +# include +# define _LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS +#endif // _CCCL_COMPILER_MSVC || _CCCL_CUDA_COMPILER_CLANG + +// NVRTC does not know the builtins. +#if defined(_CCCL_COMPILER_NVRTC) +# define _LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS +#endif // _CCCL_COMPILER_NVRTC + +_LIBCUDACXX_BEGIN_NAMESPACE_STD + +// log + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI float log(float __x) noexcept +{ +#if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::log(__x); +#else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_logf(__x); +#endif // !_CCCL_COMPILER_MSVC +} + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI float logf(float __x) noexcept +{ +#if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::log(__x); +#else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_logf(__x); +#endif // !_CCCL_COMPILER_MSVC +} + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI double log(double __x) noexcept +{ +#if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::log(__x); +#else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_log(__x); +#endif // !_CCCL_COMPILER_MSVC +} + +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI long double log(long double __x) noexcept +{ +# if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::log(__x); +# else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_logl(__x); +# endif // !_CCCL_COMPILER_MSVC +} + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI long double logl(long double __x) noexcept +{ +# if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::log(__x); +# else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_logl(__x); +# endif // !_CCCL_COMPILER_MSVC +} +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE + +#if defined(_LIBCUDACXX_HAS_NVFP16) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI __half log(__half __x) noexcept +{ + NV_IF_ELSE_TARGET(NV_PROVIDES_SM_53, (return ::hlog(__x);), ({ + float __vf = __half2float(__x); + __vf = _CUDA_VSTD::logf(__vf); + __half_raw __ret_repr = ::__float2half_rn(__vf); + + _CUDA_VSTD::uint16_t __repr = __half_raw(__x).x; + switch (__repr) + { + case 7544: + __ret_repr.x -= 1; + break; + + default:; + } + + return __ret_repr; + })) +} +#endif // _LIBCUDACXX_HAS_NVFP16 + +#if defined(_LIBCUDACXX_HAS_NVBF16) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI __nv_bfloat16 log(__nv_bfloat16 __x) noexcept +{ + NV_IF_ELSE_TARGET( + NV_IS_DEVICE, (return ::hlog(__x);), (return __float2bfloat16(_CUDA_VSTD::logf(__bfloat162float(__x)));)) +} +#endif // _LIBCUDACXX_HAS_NVBF16 + +template = 0> +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI double log(_Integer __x) noexcept +{ +#if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::log((double) __x); +#else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_log((double) __x); +#endif // !_CCCL_COMPILER_MSVC +} + +// log10 + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI float log10(float __x) noexcept +{ +#if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::log10(__x); +#else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_log10f(__x); +#endif // !_CCCL_COMPILER_MSVC +} + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI float log10f(float __x) noexcept +{ +#if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::log10(__x); +#else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_log10f(__x); +#endif // !_CCCL_COMPILER_MSVC +} + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI double log10(double __x) noexcept +{ +#if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::log10(__x); +#else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_log10(__x); +#endif // !_CCCL_COMPILER_MSVC +} + +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI long double log10(long double __x) noexcept +{ +# if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::log10(__x); +# else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_log10l(__x); +# endif // !_CCCL_COMPILER_MSVC +} + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI long double log10l(long double __x) noexcept +{ +# if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::log10(__x); +# else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_log10l(__x); +# endif // !_CCCL_COMPILER_MSVC +} +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE + +#if defined(_LIBCUDACXX_HAS_NVFP16) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI __half log10(__half __x) noexcept +{ + NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hlog10(__x);), (return __float2half(_CUDA_VSTD::log10f(__half2float(__x)));)) +} +#endif // _LIBCUDACXX_HAS_NVFP16 + +#if defined(_LIBCUDACXX_HAS_NVBF16) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI __nv_bfloat16 log10(__nv_bfloat16 __x) noexcept +{ + NV_IF_ELSE_TARGET( + NV_IS_DEVICE, (return ::hlog10(__x);), (return __float2bfloat16(_CUDA_VSTD::log10f(__bfloat162float(__x)));)) +} +#endif // _LIBCUDACXX_HAS_NVBF16 + +template = 0> +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI double log10(_Integer __x) noexcept +{ +#if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::log10((double) __x); +#else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_log10((double) __x); +#endif // !_CCCL_COMPILER_MSVC +} + +// ilogb + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI int ilogb(float __x) noexcept +{ +#if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::ilogb(__x); +#else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_ilogbf(__x); +#endif // !_CCCL_COMPILER_MSVC +} + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI int ilogbf(float __x) noexcept +{ +#if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::ilogb(__x); +#else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_ilogbf(__x); +#endif // !_CCCL_COMPILER_MSVC +} + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI int ilogb(double __x) noexcept +{ +#if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::ilogb(__x); +#else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_ilogb(__x); +#endif // !_CCCL_COMPILER_MSVC +} + +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI int ilogb(long double __x) noexcept +{ +# if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::ilogb(__x); +# else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_ilogbl(__x); +# endif // !_CCCL_COMPILER_MSVC +} + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI int ilogbl(long double __x) noexcept +{ +# if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::ilogb(__x); +# else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_ilogbl(__x); +# endif // !_CCCL_COMPILER_MSVC +} +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE + +#if defined(_LIBCUDACXX_HAS_NVFP16) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI int ilogb(__half __x) noexcept +{ + return _CUDA_VSTD::ilogbf(__half2float(__x)); +} +#endif // defined(_LIBCUDACXX_HAS_NVFP16) + +#if defined(_LIBCUDACXX_HAS_NVBF16) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI int ilogb(__nv_bfloat16 __x) noexcept +{ + return _CUDA_VSTD::ilogbf(__bfloat162float(__x)); +} +#endif // _LIBCUDACXX_HAS_NVBF16 + +template = 0> +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI int ilogb(_Integer __x) noexcept +{ +#if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::ilogb((double) __x); +#else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_ilogb((double) __x); +#endif // !_CCCL_COMPILER_MSVC +} + +// log1p + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI float log1p(float __x) noexcept +{ +#if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::log1p(__x); +#else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_log1pf(__x); +#endif // !_CCCL_COMPILER_MSVC +} + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI float log1pf(float __x) noexcept +{ +#if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::log1p(__x); +#else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_log1pf(__x); +#endif // !_CCCL_COMPILER_MSVC +} + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI double log1p(double __x) noexcept +{ +#if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::log1p(__x); +#else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_log1p(__x); +#endif // !_CCCL_COMPILER_MSVC +} + +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI long double log1p(long double __x) noexcept +{ +# if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::log1p(__x); +# else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_log1pl(__x); +# endif // !_CCCL_COMPILER_MSVC +} + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI long double log1pf(long double __x) noexcept +{ +# if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::log1p(__x); +# else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_log1pl(__x); +# endif // !_CCCL_COMPILER_MSVC +} +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE + +#if defined(_LIBCUDACXX_HAS_NVFP16) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI __half log1p(__half __x) noexcept +{ + return __float2half(_CUDA_VSTD::log1p(__half2float(__x))); +} +#endif // _LIBCUDACXX_HAS_NVFP16 + +#if defined(_LIBCUDACXX_HAS_NVBF16) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI __nv_bfloat16 log1p(__nv_bfloat16 __x) noexcept +{ + return __float2bfloat16(_CUDA_VSTD::log1p(__bfloat162float(__x))); +} +#endif // _LIBCUDACXX_HAS_NVBF16 + +template = 0> +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI double log1p(_Integer __x) noexcept +{ +#if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::log1p((double) __x); +#else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_log1p((double) __x); +#endif // !_CCCL_COMPILER_MSVC +} + +// log2 + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI float log2(float __x) noexcept +{ +#if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::log2(__x); +#else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_log2f(__x); +#endif // !_CCCL_COMPILER_MSVC +} + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI float log2f(float __x) noexcept +{ +#if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::log2(__x); +#else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_log2f(__x); +#endif // !_CCCL_COMPILER_MSVC +} + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI double log2(double __x) noexcept +{ +#if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::log2(__x); +#else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_log2(__x); +#endif // !_CCCL_COMPILER_MSVC +} + +#if defined(_LIBCUDACXX_HAS_NVFP16) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI __half log2(__half __x) noexcept +{ + NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hlog2(__x);), (return __float2half(_CUDA_VSTD::log2f(__half2float(__x)));)) +} +#endif // _LIBCUDACXX_HAS_NVFP16 + +#if defined(_LIBCUDACXX_HAS_NVBF16) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI __nv_bfloat16 log2(__nv_bfloat16 __x) noexcept +{ + NV_IF_ELSE_TARGET( + NV_IS_DEVICE, (return ::hlog2(__x);), (return __float2bfloat16(_CUDA_VSTD::log2f(__bfloat162float(__x)));)) +} +#endif // _LIBCUDACXX_HAS_NVBF16 + +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI long double log2(long double __x) noexcept +{ +# if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::log2(__x); +# else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_log2l(__x); +# endif // !_CCCL_COMPILER_MSVC +} + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI long double log2l(long double __x) noexcept +{ +# if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::log2(__x); +# else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_log2l(__x); +# endif // !_CCCL_COMPILER_MSVC +} +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE + +template = 0> +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI double log2(_Integer __x) noexcept +{ +#if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::log2((double) __x); +#else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_log2((double) __x); +#endif // !_CCCL_COMPILER_MSVC +} + +// logb + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI float logb(float __x) noexcept +{ +#if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::logb(__x); +#else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_logbf(__x); +#endif // !_CCCL_COMPILER_MSVC +} + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI float logbf(float __x) noexcept +{ +#if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::logb(__x); +#else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_logbf(__x); +#endif // !_CCCL_COMPILER_MSVC +} + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI double logb(double __x) noexcept +{ +#if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::logb(__x); +#else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_logb(__x); +#endif // !_CCCL_COMPILER_MSVC +} + +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI long double logb(long double __x) noexcept +{ +# if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::logb(__x); +# else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_logbl(__x); +# endif // !_CCCL_COMPILER_MSVC +} + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI long double logbl(long double __x) noexcept +{ +# if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::logb(__x); +# else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_logbl(__x); +# endif // !_CCCL_COMPILER_MSVC +} +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE + +#if defined(_LIBCUDACXX_HAS_NVFP16) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI __half logb(__half __x) noexcept +{ + return __float2half(_CUDA_VSTD::logbf(__half2float(__x))); +} +#endif // _LIBCUDACXX_HAS_NVFP16 + +#if defined(_LIBCUDACXX_HAS_NVBF16) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI __nv_bfloat16 logb(__nv_bfloat16 __x) noexcept +{ + return __float2bfloat16(_CUDA_VSTD::logbf(__bfloat162float(__x))); +} +#endif // _LIBCUDACXX_HAS_NVBF16 + +template = 0> +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI double logb(_Integer __x) noexcept +{ +#if defined(_LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS) + return ::logb((double) __x); +#else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + return __builtin_logb((double) __x); +#endif // !_CCCL_COMPILER_MSVC +} + +_LIBCUDACXX_END_NAMESPACE_STD + +#undef _LIBCUDACXX_CMATH_USE_CRT_FUNCTIONS + +#endif // _LIBCUDACXX___CMATH_LOGARITHMS_H diff --git a/libcudacxx/include/cuda/std/__cmath/nvbf16.h b/libcudacxx/include/cuda/std/__cmath/nvbf16.h index 66356663e43..f90fe066c4e 100644 --- a/libcudacxx/include/cuda/std/__cmath/nvbf16.h +++ b/libcudacxx/include/cuda/std/__cmath/nvbf16.h @@ -29,6 +29,7 @@ _CCCL_DIAG_SUPPRESS_CLANG("-Wunused-function") _CCCL_DIAG_POP # include +# include # include # include @@ -71,11 +72,6 @@ _LIBCUDACXX_HIDE_FROM_ABI __nv_bfloat16 atan2(__nv_bfloat16 __x, __nv_bfloat16 _ return __float2bfloat16(::atan2f(__bfloat162float(__x), __bfloat162float(__y))); } -_LIBCUDACXX_HIDE_FROM_ABI __nv_bfloat16 log(__nv_bfloat16 __x) -{ - NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hlog(__x);), (return __float2bfloat16(::logf(__bfloat162float(__x)));)) -} - _LIBCUDACXX_HIDE_FROM_ABI __nv_bfloat16 sqrt(__nv_bfloat16 __x) { NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hsqrt(__x);), (return __float2bfloat16(::sqrtf(__bfloat162float(__x)));)) diff --git a/libcudacxx/include/cuda/std/__cmath/nvfp16.h b/libcudacxx/include/cuda/std/__cmath/nvfp16.h index 2eacda3da1b..cbd48c69daf 100644 --- a/libcudacxx/include/cuda/std/__cmath/nvfp16.h +++ b/libcudacxx/include/cuda/std/__cmath/nvfp16.h @@ -26,6 +26,7 @@ # include # include +# include # include # include @@ -136,33 +137,6 @@ _LIBCUDACXX_HIDE_FROM_ABI __half atan2(__half __x, __half __y) return __float2half(::atan2f(__half2float(__x), __half2float(__y))); } -// clang-format off -_LIBCUDACXX_HIDE_FROM_ABI __half log(__half __x) -{ - NV_IF_ELSE_TARGET(NV_PROVIDES_SM_53, ( - return ::hlog(__x); - ), ( - { - float __vf = __half2float(__x); - __vf = ::logf(__vf); - __half_raw __ret_repr = ::__float2half_rn(__vf); - - uint16_t __repr = __half_raw(__x).x; - switch (__repr) - { - case 7544: - __ret_repr.x -= 1; - break; - - default:; - } - - return __ret_repr; - } - )) -} -// clang-format on - _LIBCUDACXX_HIDE_FROM_ABI __half sqrt(__half __x) { NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hsqrt(__x);), (return __float2half(::sqrtf(__half2float(__x)));)) diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/cmath b/libcudacxx/include/cuda/std/detail/libcxx/include/cmath index f33a60b8d58..5b042bae246 100644 --- a/libcudacxx/include/cuda/std/detail/libcxx/include/cmath +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/cmath @@ -387,9 +387,6 @@ using ::asinhf; using ::atanh; using ::atanhf; -using ::log; -using ::logf; - using ::hypot; using ::hypotf; @@ -425,11 +422,6 @@ using ::frexpf; using ::ldexp; using ::ldexpf; -using ::log; -using ::logf; - -using ::log10; -using ::log10f; using ::modf; using ::modff; @@ -477,20 +469,12 @@ using ::fmax; using ::fmaxf; using ::fmin; using ::fminf; -using ::ilogb; -using ::ilogbf; using ::lgamma; using ::lgammaf; using ::llrint; using ::llrintf; using ::llround; using ::llroundf; -using ::log1p; -using ::log1pf; -using ::log2; -using ::log2f; -using ::logb; -using ::logbf; using ::lrint; using ::lrintf; using ::lround; @@ -535,8 +519,6 @@ using ::floorl; using ::fmodl; using ::frexpl; using ::ldexpl; -using ::log10l; -using ::logl; using ::modfl; using ::powl; using ::sinhl; @@ -561,13 +543,9 @@ using ::fmal; using ::fmaxl; using ::fminl; using ::hypotl; -using ::ilogbl; using ::lgammal; using ::llrintl; using ::llroundl; -using ::log1pl; -using ::log2l; -using ::logbl; using ::lrintl; using ::lroundl; using ::nanl; @@ -836,7 +814,7 @@ _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14_COMPLEX __promote_t<_Tp, _Up> __ template _LIBCUDACXX_HIDE_FROM_ABI _A1 __constexpr_logb(_A1 __x) { - return ::logb(__x); + return _CUDA_VSTD::logb(__x); } #else template diff --git a/libcudacxx/test/libcudacxx/std/numerics/c.math/logarithms.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/c.math/logarithms.pass.cpp new file mode 100644 index 00000000000..356cf9386cb --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/numerics/c.math/logarithms.pass.cpp @@ -0,0 +1,205 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, 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) 2024 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// + +#include +#include +#include + +#include "test_macros.h" + +__host__ __device__ void test_log() +{ + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + static_assert((cuda::std::is_same::value), ""); +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE + static_assert((cuda::std::is_same::value), ""); +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + static_assert((cuda::std::is_same::value), ""); +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE +#ifdef _LIBCUDACXX_HAS_NVFP16 + static_assert((cuda::std::is_same::value), ""); +#endif // _LIBCUDACXX_HAS_NVFP16 +#ifdef _LIBCUDACXX_HAS_NVBF16 + static_assert((cuda::std::is_same::value), ""); +#endif // _LIBCUDACXX_HAS_NVBF16 + assert(cuda::std::log(1) == 0); +} + +__host__ __device__ void test_log10() +{ + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + static_assert((cuda::std::is_same::value), ""); +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE + static_assert((cuda::std::is_same::value), ""); +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + static_assert((cuda::std::is_same::value), ""); +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE +#ifdef _LIBCUDACXX_HAS_NVFP16 + static_assert((cuda::std::is_same::value), ""); +#endif // _LIBCUDACXX_HAS_NVFP16 +#ifdef _LIBCUDACXX_HAS_NVBF16 + static_assert((cuda::std::is_same::value), ""); +#endif // _LIBCUDACXX_HAS_NVBF16 + assert(cuda::std::log10(1) == 0); +} + +__host__ __device__ void test_ilogb() +{ + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + static_assert((cuda::std::is_same::value), ""); +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE + static_assert((cuda::std::is_same::value), ""); +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + static_assert((cuda::std::is_same::value), ""); +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE +#ifdef _LIBCUDACXX_HAS_NVFP16 + static_assert((cuda::std::is_same::value), ""); +#endif // _LIBCUDACXX_HAS_NVFP16 +#ifdef _LIBCUDACXX_HAS_NVBF16 + static_assert((cuda::std::is_same::value), ""); +#endif // _LIBCUDACXX_HAS_NVBF16 + assert(cuda::std::ilogb(1) == 0); +} + +__host__ __device__ void test_log1p() +{ + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + static_assert((cuda::std::is_same::value), ""); +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE + static_assert((cuda::std::is_same::value), ""); +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + static_assert((cuda::std::is_same::value), ""); +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE +#ifdef _LIBCUDACXX_HAS_NVFP16 + static_assert((cuda::std::is_same::value), ""); +#endif // _LIBCUDACXX_HAS_NVFP16 +#ifdef _LIBCUDACXX_HAS_NVBF16 + static_assert((cuda::std::is_same::value), ""); +#endif // _LIBCUDACXX_HAS_NVBF16 + assert(cuda::std::log1p(0) == 0); +} + +__host__ __device__ void test_log2() +{ + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + static_assert((cuda::std::is_same::value), ""); +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE + static_assert((cuda::std::is_same::value), ""); +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + static_assert((cuda::std::is_same::value), ""); +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE +#ifdef _LIBCUDACXX_HAS_NVFP16 + static_assert((cuda::std::is_same::value), ""); +#endif // _LIBCUDACXX_HAS_NVFP16 +#ifdef _LIBCUDACXX_HAS_NVBF16 + static_assert((cuda::std::is_same::value), ""); +#endif // _LIBCUDACXX_HAS_NVBF16 + assert(cuda::std::log2(1) == 0); +} + +__host__ __device__ void test_logb() +{ + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + static_assert((cuda::std::is_same::value), ""); +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE + static_assert((cuda::std::is_same::value), ""); +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + static_assert((cuda::std::is_same::value), ""); +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE +#ifdef _LIBCUDACXX_HAS_NVFP16 + static_assert((cuda::std::is_same::value), ""); +#endif // _LIBCUDACXX_HAS_NVFP16 +#ifdef _LIBCUDACXX_HAS_NVBF16 + static_assert((cuda::std::is_same::value), ""); +#endif // _LIBCUDACXX_HAS_NVBF16 + assert(cuda::std::logb(1) == 0); +} + +__host__ __device__ void test() +{ + test_log(); + test_log10(); + test_ilogb(); + test_log1p(); + test_log2(); + test_logb(); +} + +__global__ void test_global_kernel() +{ + test(); +} + +int main(int, char**) +{ + test(); + return 0; +}