-
Notifications
You must be signed in to change notification settings - Fork 184
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
Implement more cmath functions to be usable on host and device (#3382)
* Implement more cmath functions to be usable on host and device * Implement math roots functions * Implement exponential functions
- Loading branch information
Showing
10 changed files
with
3,103 additions
and
131 deletions.
There are no files selected for viewing
Large diffs are not rendered by default.
Oops, something went wrong.
611 changes: 611 additions & 0 deletions
611
libcudacxx/include/cuda/std/__cmath/exponential_functions.h
Large diffs are not rendered by default.
Oops, something went wrong.
This file contains 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
This file contains 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
This file contains 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
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,171 @@ | ||
| // -*- 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) 2025 NVIDIA CORPORATION & AFFILIATES. | ||
| // | ||
| //===----------------------------------------------------------------------===// | ||
|
|
||
| #ifndef _LIBCUDACXX___CMATH_ROOTS_H | ||
| #define _LIBCUDACXX___CMATH_ROOTS_H | ||
|
|
||
| #include <cuda/std/detail/__config> | ||
|
|
||
| #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 <cuda/std/__cmath/common.h> | ||
| #include <cuda/std/__type_traits/enable_if.h> | ||
| #include <cuda/std/__type_traits/is_integral.h> | ||
|
|
||
| _LIBCUDACXX_BEGIN_NAMESPACE_STD | ||
|
|
||
| // sqrt | ||
|
|
||
| _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI float sqrt(float __x) noexcept | ||
| { | ||
| #if defined(_CCCL_BUILTIN_SQRTF) | ||
| return _CCCL_BUILTIN_SQRTF(__x); | ||
| #else // ^^^ _CCCL_BUILTIN_SQRTF ^^^ // vvv !_CCCL_BUILTIN_SQRTF vvv | ||
| return ::sqrtf(__x); | ||
| #endif // !_CCCL_BUILTIN_SQRTF | ||
| } | ||
|
|
||
| _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI float sqrtf(float __x) noexcept | ||
| { | ||
| #if defined(_CCCL_BUILTIN_SQRTF) | ||
| return _CCCL_BUILTIN_SQRTF(__x); | ||
| #else // ^^^ _CCCL_BUILTIN_SQRTF ^^^ // vvv !_CCCL_BUILTIN_SQRTF vvv | ||
| return ::sqrtf(__x); | ||
| #endif // !_CCCL_BUILTIN_SQRTF | ||
| } | ||
|
|
||
| _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI double sqrt(double __x) noexcept | ||
| { | ||
| #if defined(_CCCL_BUILTIN_SQRT) | ||
| return _CCCL_BUILTIN_SQRT(__x); | ||
| #else // ^^^ _CCCL_BUILTIN_SQRT ^^^ // vvv !_CCCL_BUILTIN_SQRT vvv | ||
| return ::sqrt(__x); | ||
| #endif // !_CCCL_BUILTIN_SQRT | ||
| } | ||
|
|
||
| #if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) | ||
| _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI long double sqrt(long double __x) noexcept | ||
| { | ||
| # if defined(_CCCL_BUILTIN_SQRTL) | ||
| return _CCCL_BUILTIN_SQRTL(__x); | ||
| # else // ^^^ _CCCL_BUILTIN_SQRTL ^^^ // vvv !_CCCL_BUILTIN_SQRTL vvv | ||
| return ::sqrtl(__x); | ||
| # endif // !_CCCL_BUILTIN_SQRTL | ||
| } | ||
|
|
||
| _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI long double sqrtl(long double __x) noexcept | ||
| { | ||
| # if defined(_CCCL_BUILTIN_SQRTL) | ||
| return _CCCL_BUILTIN_SQRTL(__x); | ||
| # else // ^^^ _CCCL_BUILTIN_SQRTL ^^^ // vvv !_CCCL_BUILTIN_SQRTL vvv | ||
| return ::sqrtl(__x); | ||
| # endif // !_CCCL_BUILTIN_SQRTL | ||
| } | ||
| #endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE | ||
|
|
||
| #if defined(_LIBCUDACXX_HAS_NVFP16) | ||
| _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI __half sqrt(__half __x) noexcept | ||
| { | ||
| NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hsqrt(__x);), (return __float2half(_CUDA_VSTD::sqrt(__half2float(__x)));)) | ||
| } | ||
| #endif // _LIBCUDACXX_HAS_NVFP16 | ||
|
|
||
| #if defined(_LIBCUDACXX_HAS_NVBF16) | ||
| _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI __nv_bfloat16 sqrt(__nv_bfloat16 __x) noexcept | ||
| { | ||
| NV_IF_ELSE_TARGET( | ||
| NV_IS_DEVICE, (return ::hsqrt(__x);), (return __float2bfloat16(_CUDA_VSTD::sqrt(__bfloat162float(__x)));)) | ||
| } | ||
| #endif // _LIBCUDACXX_HAS_NVBF16 | ||
|
|
||
| template <class _Integer, enable_if_t<_CCCL_TRAIT(is_integral, _Integer), int> = 0> | ||
| _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI double sqrt(_Integer __x) noexcept | ||
| { | ||
| return _CUDA_VSTD::sqrt((double) __x); | ||
| } | ||
|
|
||
| // cbrt | ||
|
|
||
| _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI float cbrt(float __x) noexcept | ||
| { | ||
| #if defined(_CCCL_BUILTIN_CBRTF) | ||
| return _CCCL_BUILTIN_CBRTF(__x); | ||
| #else // ^^^ _CCCL_BUILTIN_CBRTF ^^^ // vvv !_CCCL_BUILTIN_CBRTF vvv | ||
| return ::cbrtf(__x); | ||
| #endif // !_CCCL_BUILTIN_CBRTF | ||
| } | ||
|
|
||
| _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI float cbrtf(float __x) noexcept | ||
| { | ||
| #if defined(_CCCL_BUILTIN_CBRTF) | ||
| return _CCCL_BUILTIN_CBRTF(__x); | ||
| #else // ^^^ _CCCL_BUILTIN_CBRTF ^^^ // vvv !_CCCL_BUILTIN_CBRTF vvv | ||
| return ::cbrtf(__x); | ||
| #endif // !_CCCL_BUILTIN_CBRTF | ||
| } | ||
|
|
||
| _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI double cbrt(double __x) noexcept | ||
| { | ||
| #if defined(_CCCL_BUILTIN_CBRT) | ||
| return _CCCL_BUILTIN_CBRT(__x); | ||
| #else // ^^^ _CCCL_BUILTIN_CBRT ^^^ // vvv !_CCCL_BUILTIN_CBRT vvv | ||
| return ::cbrt(__x); | ||
| #endif // !_CCCL_BUILTIN_CBRT | ||
| } | ||
|
|
||
| #if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) | ||
| _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI long double cbrt(long double __x) noexcept | ||
| { | ||
| # if defined(_CCCL_BUILTIN_CBRTL) | ||
| return _CCCL_BUILTIN_CBRTL(__x); | ||
| # else // ^^^ _CCCL_BUILTIN_CBRTL ^^^ // vvv !_CCCL_BUILTIN_CBRTL vvv | ||
| return ::cbrtl(__x); | ||
| # endif // !_CCCL_BUILTIN_CBRTL | ||
| } | ||
|
|
||
| _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI long double cbrtl(long double __x) noexcept | ||
| { | ||
| # if defined(_CCCL_BUILTIN_CBRTL) | ||
| return _CCCL_BUILTIN_CBRTL(__x); | ||
| # else // ^^^ _CCCL_BUILTIN_CBRTL ^^^ // vvv !_CCCL_BUILTIN_CBRTL vvv | ||
| return ::cbrtl(__x); | ||
| # endif // !_CCCL_BUILTIN_CBRTL | ||
| } | ||
| #endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE | ||
|
|
||
| #if defined(_LIBCUDACXX_HAS_NVFP16) | ||
| _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI __half cbrt(__half __x) noexcept | ||
| { | ||
| return __float2half(_CUDA_VSTD::cbrt(__half2float(__x))); | ||
| } | ||
| #endif // _LIBCUDACXX_HAS_NVFP16 | ||
|
|
||
| #if defined(_LIBCUDACXX_HAS_NVBF16) | ||
| _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI __nv_bfloat16 cbrt(__nv_bfloat16 __x) noexcept | ||
| { | ||
| return __float2bfloat16(_CUDA_VSTD::cbrt(__bfloat162float(__x))); | ||
| } | ||
| #endif // _LIBCUDACXX_HAS_NVBF16 | ||
|
|
||
| template <class _Integer, enable_if_t<_CCCL_TRAIT(is_integral, _Integer), int> = 0> | ||
| _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI double cbrt(_Integer __x) noexcept | ||
| { | ||
| return _CUDA_VSTD::cbrt((double) __x); | ||
| } | ||
|
|
||
| _LIBCUDACXX_END_NAMESPACE_STD | ||
|
|
||
| #endif // _LIBCUDACXX___CMATH_ROOTS_H |
Oops, something went wrong.