1#ifndef TDLS_SOLVERS_TILE_OPERATIONS_HPP
2#define TDLS_SOLVERS_TILE_OPERATIONS_HPP
39#pragma clang diagnostic push
40#pragma clang diagnostic ignored "-Wpass-failed"
52template<
typename T,
int tile_size,
bool unroll_inner>
66 template<
int k_extent>
69 if constexpr (unroll_inner) {
71 for (
int row = 0; row < k_extent; ++row) {
74 for (
int j = 0; j < k_extent; ++j) {
75 const T tmp = t[k * tile_size + j];
76 t[k * tile_size + j] = t[row * tile_size + j];
77 t[row * tile_size + j] = tmp;
82 for (
int row = 0; row < k_extent; ++row) {
84 for (
int j = 0; j < k_extent; ++j) {
85 const T tmp = t[k * tile_size + j];
86 t[k * tile_size + j] = t[row * tile_size + j];
87 t[row * tile_size + j] = tmp;
100 template<
int diag_extent,
int col_extent>
103 if constexpr (unroll_inner) {
105 for (
int k = 0; k < diag_extent; ++k) {
107 for (
int i = k + 1; i < diag_extent; ++i) {
108 const T L_ik = lu[i * tile_size + k];
110 for (
int j = 0; j < col_extent; ++j)
111 B[i * tile_size + j] -= L_ik * B[k * tile_size + j];
115 for (
int k = 0; k < diag_extent; ++k) {
116 for (
int i = k + 1; i < diag_extent; ++i) {
117 const T L_ik = lu[i * tile_size + k];
118 for (
int j = 0; j < col_extent; ++j)
119 B[i * tile_size + j] -= L_ik * B[k * tile_size + j];
134 template<
int diag_extent,
int row_extent>
137 if constexpr (unroll_inner) {
139 for (
int k = 0; k < diag_extent; ++k) {
140 const T U_kk_inv = lu[k * tile_size + k];
142 for (
int i = 0; i < row_extent; ++i)
143 B[i * tile_size + k] *= U_kk_inv;
145 for (
int j = k + 1; j < diag_extent; ++j) {
146 const T U_kj = lu[k * tile_size + j];
148 for (
int i = 0; i < row_extent; ++i)
149 B[i * tile_size + j] -= B[i * tile_size + k] * U_kj;
153 for (
int k = 0; k < diag_extent; ++k) {
154 const T U_kk_inv = lu[k * tile_size + k];
155 for (
int i = 0; i < row_extent; ++i)
156 B[i * tile_size + k] *= U_kk_inv;
157 for (
int j = k + 1; j < diag_extent; ++j) {
158 const T U_kj = lu[k * tile_size + j];
159 for (
int i = 0; i < row_extent; ++i)
160 B[i * tile_size + j] -= B[i * tile_size + k] * U_kj;
174 template<
int row_extent,
int col_extent,
int k_extent>
177 if constexpr (unroll_inner) {
179 for (
int i = 0; i < row_extent; ++i) {
181 for (
int j = 0; j < col_extent; ++j) {
184 for (
int k = 0; k < k_extent; ++k)
185 sum += At[i * tile_size + k] * Bt[k * tile_size + j];
186 Ct[i * tile_size + j] -= sum;
190 for (
int i = 0; i < row_extent; ++i) {
191 for (
int j = 0; j < col_extent; ++j) {
193 for (
int k = 0; k < k_extent; ++k)
194 sum += At[i * tile_size + k] * Bt[k * tile_size + j];
195 Ct[i * tile_size + j] -= sum;
206#if defined(__clang__)
207#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
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void gemm_sub(T *TDLS_RESTRICT Ct, const T *TDLS_RESTRICT At, const T *TDLS_RESTRICT Bt) noexcept
Ct -= At*Bt with per-element dot-product accumulation. At is RxK, Bt is KxC, Ct is RxC.
Definition tile_operations.hpp:176
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void swap_rows(T *TDLS_RESTRICT t, int k, int r) noexcept
Row swap k <-> r inside the KExKE active part of the tile, compile-time indexed on both sides.
Definition tile_operations.hpp:67
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void trsm_right(const T *TDLS_RESTRICT lu, T *TDLS_RESTRICT B) noexcept
B := B U^-1, with U the upper part of the factored diagonal tile. U is KDxKD, B is RxKD.
Definition tile_operations.hpp:136
TDLS_HOST_DEVICE static TDLS_FORCEINLINE constexpr void trsm_left_unit(const T *TDLS_RESTRICT lu, T *TDLS_RESTRICT B) noexcept
B := L^-1 B, with L the unit lower part of the factored diagonal tile. L is KDxKD,...
Definition tile_operations.hpp:102