TDLS 0.4.0
Tiny Device-callable Linear Solvers
Loading...
Searching...
No Matches
options.hpp
Go to the documentation of this file.
1#ifndef TDLS_SOLVERS_OPTIONS_HPP
2#define TDLS_SOLVERS_OPTIONS_HPP
3
4
5
32
33
34
35#include <type_traits>
36
37#include <tdls/core/macros.hpp>
38#include <tdls/core/target.hpp>
39
40#if defined(TDLS_DETAIL_SIDE_DISPATCH)
41#include <nv/target>
42#endif
43
44
45
46namespace tdls {
47
48
49
51enum class Schedule {
54};
55
56
57
59enum class RowInterchange {
60 Logical,
61 Physical,
67 Auto
68};
69
70namespace detail {
71
83template<typename T, bool gpu>
84[[nodiscard]] TDLS_HOST_DEVICE constexpr bool physical_rows_fit_on(const int n) noexcept {
85 if constexpr (gpu)
86 return sizeof(T) <= 4 ? n <= 14 : n <= 10;
87 else
88 return sizeof(T) <= 4 ? n <= 8 : n <= 12;
89}
90
96template<typename T>
97[[nodiscard]] TDLS_HOST_DEVICE constexpr bool physical_rows_fit(const int n) noexcept {
98 return physical_rows_fit_on<T, TDLS_TARGET_GPU != 0>(n);
99}
100
108template<typename T, bool gpu>
109[[nodiscard]] TDLS_HOST_DEVICE constexpr bool physical_rows_on(const RowInterchange r,
110 const int n) noexcept {
111 return r == RowInterchange::Physical ||
112 (r == RowInterchange::Auto && physical_rows_fit_on<T, gpu>(n));
113}
114
121template<typename T>
122[[nodiscard]] TDLS_HOST_DEVICE constexpr bool physical_rows(const RowInterchange r,
123 const int n) noexcept {
124 return physical_rows_on<T, TDLS_TARGET_GPU != 0>(r, n);
125}
126
127#if defined(TDLS_DETAIL_SIDE_DISPATCH)
133TDLS_HOST_DEVICE inline bool on_device_side() noexcept {
134 NV_IF_TARGET(NV_IS_DEVICE, (return true;), (return false;))
135}
136#endif
137
146template<typename T>
147[[nodiscard]] TDLS_HOST_DEVICE constexpr bool physical_rows_here(const RowInterchange r,
148 const int n) noexcept {
149#if defined(TDLS_DETAIL_SIDE_DISPATCH)
150 if (!std::is_constant_evaluated())
151 return on_device_side() ? physical_rows_on<T, true>(r, n)
152 : physical_rows_on<T, false>(r, n);
153#endif
154 return physical_rows<T>(r, n);
155}
156
157} // namespace detail
158
159
160
181inline constexpr int dynamic_stride = -1;
182
183namespace detail {
184
190template<int s>
191[[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_ALL constexpr int
192stride_or(const int runtime) noexcept {
193 static_assert(s == dynamic_stride || s >= 1,
194 "tdls: a compile-time stride must be tdls::dynamic_stride or positive");
195 if constexpr (s == dynamic_stride)
196 return runtime;
197 else
198 return s;
199}
200
205template<int... s>
206inline constexpr bool any_static_stride = ((s != dynamic_stride) || ...);
207
208} // namespace detail
209
210
211
213enum class MatrixLayout {
214 RowMajor,
215 ColMajor
216};
217
218
219
220} // namespace tdls
221
222
223
224#endif // TDLS_SOLVERS_OPTIONS_HPP
Toolchain detection and portability macros.
#define TDLS_FORCE_INLINE_ALL
Inlining forced on every backend, CPU included: TDLS_FORCE_INLINE_GPU on GPU, __forceinline under MSV...
Definition macros.hpp:101
#define TDLS_HOST_DEVICE
__host__ __device__ under CUDA, the equivalent attributes under HIP, empty elsewhere.
Definition macros.hpp:50
TDLS_HOST_DEVICE constexpr bool physical_rows_fit(const int n) noexcept
physical_rows_fit_on for the hardware of this compilation pass (TDLS_TARGET_GPU).
Definition options.hpp:97
TDLS_HOST_DEVICE constexpr bool physical_rows_fit_on(const int n) noexcept
Whether a matrix of n unknowns is small enough for the physical row interchanges on the given hardwar...
Definition options.hpp:84
RowInterchange
Row interchanges of a factorization with partial pivoting.
Definition options.hpp:59
@ Logical
Rows never move: the position of each row in the pivoted order is tracked.
Schedule
Elimination schedule of a factorization.
Definition options.hpp:51
@ RightLooking
Factor the diagonal block, push updates into the trailing matrix.
@ LeftLooking
Pull updates from prior blocks when a block is visited.
MatrixLayout
Memory layout of a dense matrix.
Definition options.hpp:213
@ ColMajor
Element (r, c) at flat index c * N + r.
@ RowMajor
Element (r, c) at flat index r * N + c, the TFEL convention.
TDLS_HOST_DEVICE constexpr bool physical_rows_on(const RowInterchange r, const int n) noexcept
Whether a factorization of n unknowns interchanges its rows physically on the given hardware,...
Definition options.hpp:109
constexpr bool any_static_stride
Whether one of the strides of an entry point is known at compile time: the entry point then runs its ...
Definition options.hpp:206
TDLS_HOST_DEVICE constexpr bool physical_rows_here(const RowInterchange r, const int n) noexcept
Whether a factorization of n unknowns interchanges its rows physically, decided at run time on the si...
Definition options.hpp:147
TDLS_HOST_DEVICE constexpr bool physical_rows(const RowInterchange r, const int n) noexcept
physical_rows_on for the hardware of this compilation pass (TDLS_TARGET_GPU).
Definition options.hpp:122
TDLS_HOST_DEVICE TDLS_FORCE_INLINE_ALL constexpr int stride_or(const int runtime) noexcept
The stride an entry point uses: the compile-time stride when it is known, the runtime one otherwise (...
Definition options.hpp:192
constexpr int dynamic_stride
The compile-time stride of an operand whose stride is known at run time only: the default of the stri...
Definition options.hpp:181
Detection of the compiler, of the language and of the side of the compilation pass: the one place whe...