TDLS 0.4.0
Tiny Device-callable Linear Solvers
Loading...
Searching...
No Matches
math.hpp
Go to the documentation of this file.
1#ifndef TDLS_CORE_MATH_HPP
2#define TDLS_CORE_MATH_HPP
3
4
5
13
14
15
16#include <cmath>
17
18#include <tdls/core/macros.hpp>
19#include <tdls/core/target.hpp>
20
21
22
23namespace tdls {
24
25
26
27namespace detail {
28
41template<typename T>
42[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr T abs(const T x) noexcept {
43#if defined(__has_builtin)
44#if __has_builtin(__builtin_is_constant_evaluated)
45 if (__builtin_is_constant_evaluated()) return x < T(0) ? -x : x;
46#endif
47#elif defined(TDLS_DETAIL_COMPILER_MSVC) && _MSC_VER >= 1925
48 if (__builtin_is_constant_evaluated()) return x < T(0) ? -x : x;
49#endif
50 return std::fabs(x);
51}
52
69#if defined(TDLS_DETAIL_COMPILER_NVCC) && defined(TDLS_DETAIL_PASS_CUDA) && __CUDA_ARCH__ >= 1000
70 if (!__builtin_is_constant_evaluated()) asm volatile("" : "+r"(v));
71#endif
72 return v;
73}
74
75} // namespace detail
76
77
78
79} // namespace tdls
80
81
82
83#endif // TDLS_CORE_MATH_HPP
Toolchain detection and portability macros.
#define TDLS_FORCE_INLINE_GPU
Inlining forced on GPU only: __forceinline__ under CUDA, the equivalent attribute under HIP (see TDLS...
Definition macros.hpp:77
#define TDLS_HOST_DEVICE
__host__ __device__ under CUDA, the equivalent attributes under HIP, empty elsewhere.
Definition macros.hpp:50
TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr T abs(const T x) noexcept
Absolute value usable in constant expressions.
Definition math.hpp:42
TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU constexpr int runtime_dimension(int v) noexcept
A dimension hidden from the optimizer of nvcc when it compiles device code for Blackwell (sm_100 and ...
Definition math.hpp:68
Detection of the compiler, of the language and of the side of the compilation pass: the one place whe...