TDLS 0.4.0
Tiny Device-callable Linear Solvers
Loading...
Searching...
No Matches
tile_operations.hpp
Go to the documentation of this file.
1#ifndef TDLS_SOLVERS_TILED_LUPP_TILE_OPERATIONS_HPP
2#define TDLS_SOLVERS_TILED_LUPP_TILE_OPERATIONS_HPP
3
4
5
19
20
21
22#include <tdls/core/macros.hpp>
23#include <tdls/core/target.hpp>
25
26
27
28// clang reports a forced unrolling that the optimizer could not perform
29// through -Wpass-failed. The unrolling requested by TDLS_UNROLL_FORCE is
30// a performance hint: a failed hint does not affect correctness. The
31// suppression is scoped to this header and to that warning only.
32#if defined(TDLS_DETAIL_COMPILER_CLANG)
33#pragma clang diagnostic push
34#pragma clang diagnostic ignored "-Wpass-failed"
35#endif
36
37namespace tdls {
38
39
40
46template<typename T, int tile_size, bool unroll_loops>
47struct TiledLUppTileOperations : TileOperations<T, tile_size, unroll_loops> {
48
61 template<int row_extent, int col_extent>
62 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
63 eliminate_column(T* TDLS_RESTRICT t, int k) noexcept {
64 const T inv_pivot = T(1) / t[k * tile_size + k];
65 t[k * tile_size + k] = inv_pivot;
66 if constexpr (unroll_loops) {
68 for (int i = k + 1; i < row_extent; ++i) {
69 t[i * tile_size + k] *= inv_pivot;
71 for (int j = k + 1; j < col_extent; ++j)
72 t[i * tile_size + j] -= t[i * tile_size + k] * t[k * tile_size + j];
73 }
74 } else {
75 for (int i = k + 1; i < row_extent; ++i) {
76 t[i * tile_size + k] *= inv_pivot;
77 for (int j = k + 1; j < col_extent; ++j)
78 t[i * tile_size + j] -= t[i * tile_size + k] * t[k * tile_size + j];
79 }
80 }
81 }
82};
83
84
85
86} // namespace tdls
87
88#if defined(TDLS_DETAIL_COMPILER_CLANG)
89#pragma clang diagnostic pop
90#endif
91
92
93
94#endif // TDLS_SOLVERS_TILED_LUPP_TILE_OPERATIONS_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_RESTRICT
Non-aliasing pointer qualifier (__restrict__; __restrict on MSVC; nothing in the code that nvc++ gene...
Definition macros.hpp:174
#define TDLS_UNROLL_FORCE
Full-unroll pragma in the local compiler dialect.
Definition macros.hpp:250
#define TDLS_HOST_DEVICE
__host__ __device__ under CUDA, the equivalent attributes under HIP, empty elsewhere.
Definition macros.hpp:50
tile_size x tile_size register-tile micro-kernels shared by the solver families.
Definition tile_operations.hpp:54
tile_size x tile_size register-tile micro-kernels of the TiledLUpp solvers: the shared kernels plus t...
Definition tile_operations.hpp:47
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr void eliminate_column(T *TDLS_RESTRICT t, int k) noexcept
Gaussian elimination of column k inside the diagonal tile.
Definition tile_operations.hpp:63
Detection of the compiler, of the language and of the side of the compilation pass: the one place whe...
Register-tile micro-kernels shared by the solver families.