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_TILE_OPERATIONS_HPP
2#define TDLS_SOLVERS_TILE_OPERATIONS_HPP
3
4
5
27
28
29
30#include <tdls/core/macros.hpp>
31
32
33
34// clang reports a forced unrolling that the optimizer could not perform
35// through -Wpass-failed. The unrolling requested by TDLS_UNROLL_FORCE is
36// a performance hint: a failed hint does not affect correctness. The
37// suppression is scoped to this header and to that warning only.
38#if defined(__clang__)
39#pragma clang diagnostic push
40#pragma clang diagnostic ignored "-Wpass-failed"
41#endif
42
43namespace tdls {
44
45
46
52template<typename T, int tile_size, bool unroll_inner>
54
66 template<int k_extent>
67 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void swap_rows(T* TDLS_RESTRICT t, int k,
68 int r) noexcept {
69 if constexpr (unroll_inner) {
71 for (int row = 0; row < k_extent; ++row) {
72 if (row == r) {
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;
78 }
79 }
80 }
81 } else {
82 for (int row = 0; row < k_extent; ++row) {
83 if (row == r) {
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;
88 }
89 }
90 }
91 }
92 }
93
100 template<int diag_extent, int col_extent>
101 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
102 trsm_left_unit(const T* TDLS_RESTRICT lu, T* TDLS_RESTRICT B) noexcept {
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];
112 }
113 }
114 } else {
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];
120 }
121 }
122 }
123 }
124
134 template<int diag_extent, int row_extent>
135 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
136 trsm_right(const T* TDLS_RESTRICT lu, T* TDLS_RESTRICT B) noexcept {
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;
150 }
151 }
152 } else {
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;
161 }
162 }
163 }
164 }
165
174 template<int row_extent, int col_extent, int k_extent>
175 TDLS_HOST_DEVICE TDLS_FORCEINLINE static constexpr void
176 gemm_sub(T* TDLS_RESTRICT Ct, const T* TDLS_RESTRICT At, const T* TDLS_RESTRICT Bt) noexcept {
177 if constexpr (unroll_inner) {
179 for (int i = 0; i < row_extent; ++i) {
181 for (int j = 0; j < col_extent; ++j) {
182 T sum = T(0);
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;
187 }
188 }
189 } else {
190 for (int i = 0; i < row_extent; ++i) {
191 for (int j = 0; j < col_extent; ++j) {
192 T sum = T(0);
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;
196 }
197 }
198 }
199 }
200};
201
202
203
204} // namespace tdls
205
206#if defined(__clang__)
207#pragma clang diagnostic pop
208#endif
209
210
211
212#endif // TDLS_SOLVERS_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
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