1#ifndef TDLS_SOLVERS_TILED_LUPP_TILE_OPERATIONS_HPP
2#define TDLS_SOLVERS_TILED_LUPP_TILE_OPERATIONS_HPP
32#pragma clang diagnostic push
33#pragma clang diagnostic ignored "-Wpass-failed"
45template<
typename T,
int tile_size,
bool unroll_inner>
60 template<
int row_extent,
int col_extent>
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];
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];
88#pragma clang diagnostic pop
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.