1#ifndef TDLS_SOLVERS_TILED_LUPP_TILE_OPERATIONS_HPP
2#define TDLS_SOLVERS_TILED_LUPP_TILE_OPERATIONS_HPP
32#if defined(TDLS_DETAIL_COMPILER_CLANG)
33#pragma clang diagnostic push
34#pragma clang diagnostic ignored "-Wpass-failed"
46template<
typename T,
int tile_size,
bool unroll_loops>
61 template<
int row_extent,
int col_extent>
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];
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];
88#if defined(TDLS_DETAIL_COMPILER_CLANG)
89#pragma clang diagnostic pop
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.