TDLS 0.2.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>
24
25
26
27// clang reports a forced unrolling that the optimizer could not perform
28// through -Wpass-failed. The unrolling requested by TDLS_UNROLL_FORCE is
29// a performance hint: a failed hint does not affect correctness. The
30// suppression is scoped to this header and to that warning only.
31#if defined(__clang__)
32#pragma clang diagnostic push
33#pragma clang diagnostic ignored "-Wpass-failed"
34#endif
35
36namespace tdls {
37
38
39
45template<typename T, int tile_size, bool unroll_inner>
46struct TiledLUppTileOperations : TileOperations<T, tile_size, unroll_inner> {
47
60 template<int row_extent, int col_extent>
62 int k) noexcept {
63 const T inv_pivot = T(1) / t[k * tile_size + k];
64 t[k * tile_size + k] = inv_pivot;
65 if constexpr (unroll_inner) {
67 for (int i = k + 1; i < row_extent; ++i) {
68 t[i * tile_size + k] *= inv_pivot;
70 for (int j = k + 1; j < col_extent; ++j)
71 t[i * tile_size + j] -= t[i * tile_size + k] * t[k * tile_size + j];
72 }
73 } else {
74 for (int i = k + 1; i < row_extent; ++i) {
75 t[i * tile_size + k] *= inv_pivot;
76 for (int j = k + 1; j < col_extent; ++j)
77 t[i * tile_size + j] -= t[i * tile_size + k] * t[k * tile_size + j];
78 }
79 }
80 }
81};
82
83
84
85} // namespace tdls
86
87#if defined(__clang__)
88#pragma clang diagnostic pop
89#endif
90
91
92
93#endif // TDLS_SOLVERS_TILED_LUPP_TILE_OPERATIONS_HPP
Toolchain detection and portability macros.
#define TDLS_RESTRICT
Non-aliasing pointer qualifier (__restrict__; __restrict on MSVC).
Definition macros.hpp:80
#define TDLS_UNROLL_FORCE
Full-unroll pragma in the local compiler dialect.
Definition macros.hpp:103
#define TDLS_HOST_DEVICE
__host__ __device__ under CUDA, the equivalent attributes under HIP, empty elsewhere.
Definition macros.hpp:45
#define TDLS_FORCEINLINE
__forceinline__ under CUDA, the equivalent attribute under HIP (see TDLS_HOST_DEVICE for why),...
Definition macros.hpp:67
tile_size x tile_size register-tile micro-kernels shared by the solver families.
Definition tile_operations.hpp:53
tile_size x tile_size register-tile micro-kernels of the TiledLUpp solvers: the shared kernels plus t...
Definition tile_operations.hpp:46
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void eliminate_column(T *TDLS_RESTRICT t, int k) noexcept
Gaussian elimination of column k inside the diagonal tile.
Definition tile_operations.hpp:61
Register-tile micro-kernels shared by the solver families.