TDLS 0.4.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#include <tdls/core/target.hpp>
32
33
34
35// clang reports a forced unrolling that the optimizer could not perform
36// through -Wpass-failed. The unrolling requested by TDLS_UNROLL_FORCE is
37// a performance hint: a failed hint does not affect correctness. The
38// suppression is scoped to this header and to that warning only.
39#if defined(TDLS_DETAIL_COMPILER_CLANG)
40#pragma clang diagnostic push
41#pragma clang diagnostic ignored "-Wpass-failed"
42#endif
43
44namespace tdls {
45
46
47
53template<typename T, int tile_size, bool unroll_loops>
55
67 template<int k_extent>
69 int k, int r) noexcept {
70 if constexpr (unroll_loops) {
72 for (int row = 0; row < k_extent; ++row) {
73 if (row == r) {
75 for (int j = 0; j < k_extent; ++j) {
76 const T tmp = t[k * tile_size + j];
77 t[k * tile_size + j] = t[row * tile_size + j];
78 t[row * tile_size + j] = tmp;
79 }
80 }
81 }
82 } else {
83 for (int row = 0; row < k_extent; ++row) {
84 if (row == r) {
85 for (int j = 0; j < k_extent; ++j) {
86 const T tmp = t[k * tile_size + j];
87 t[k * tile_size + j] = t[row * tile_size + j];
88 t[row * tile_size + j] = tmp;
89 }
90 }
91 }
92 }
93 }
94
101 template<int diag_extent, int col_extent>
102 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
103 trsm_left_unit(const T* TDLS_RESTRICT lu, T* TDLS_RESTRICT B) noexcept {
104 if constexpr (unroll_loops) {
106 for (int k = 0; k < diag_extent; ++k) {
108 for (int i = k + 1; i < diag_extent; ++i) {
109 const T L_ik = lu[i * tile_size + k];
111 for (int j = 0; j < col_extent; ++j)
112 B[i * tile_size + j] -= L_ik * B[k * tile_size + j];
113 }
114 }
115 } else {
116 for (int k = 0; k < diag_extent; ++k) {
117 for (int i = k + 1; i < diag_extent; ++i) {
118 const T L_ik = lu[i * tile_size + k];
119 for (int j = 0; j < col_extent; ++j)
120 B[i * tile_size + j] -= L_ik * B[k * tile_size + j];
121 }
122 }
123 }
124 }
125
135 template<int diag_extent, int row_extent>
136 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
137 trsm_right(const T* TDLS_RESTRICT lu, T* TDLS_RESTRICT B) noexcept {
138 if constexpr (unroll_loops) {
140 for (int k = 0; k < diag_extent; ++k) {
141 const T U_kk_inv = lu[k * tile_size + k];
143 for (int i = 0; i < row_extent; ++i)
144 B[i * tile_size + k] *= U_kk_inv;
146 for (int j = k + 1; j < diag_extent; ++j) {
147 const T U_kj = lu[k * tile_size + j];
149 for (int i = 0; i < row_extent; ++i)
150 B[i * tile_size + j] -= B[i * tile_size + k] * U_kj;
151 }
152 }
153 } else {
154 for (int k = 0; k < diag_extent; ++k) {
155 const T U_kk_inv = lu[k * tile_size + k];
156 for (int i = 0; i < row_extent; ++i)
157 B[i * tile_size + k] *= U_kk_inv;
158 for (int j = k + 1; j < diag_extent; ++j) {
159 const T U_kj = lu[k * tile_size + j];
160 for (int i = 0; i < row_extent; ++i)
161 B[i * tile_size + j] -= B[i * tile_size + k] * U_kj;
162 }
163 }
164 }
165 }
166
175 template<int row_extent, int col_extent, int k_extent>
176 TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr void
177 gemm_sub(T* TDLS_RESTRICT Ct, const T* TDLS_RESTRICT At, const T* TDLS_RESTRICT Bt) noexcept {
178 if constexpr (unroll_loops) {
180 for (int i = 0; i < row_extent; ++i) {
182 for (int j = 0; j < col_extent; ++j) {
183 T sum = T(0);
185 for (int k = 0; k < k_extent; ++k)
186 sum += At[i * tile_size + k] * Bt[k * tile_size + j];
187 Ct[i * tile_size + j] -= sum;
188 }
189 }
190 } else {
191 for (int i = 0; i < row_extent; ++i) {
192 for (int j = 0; j < col_extent; ++j) {
193 T sum = T(0);
194 for (int k = 0; k < k_extent; ++k)
195 sum += At[i * tile_size + k] * Bt[k * tile_size + j];
196 Ct[i * tile_size + j] -= sum;
197 }
198 }
199 }
200 }
201};
202
203
204
205} // namespace tdls
206
207#if defined(TDLS_DETAIL_COMPILER_CLANG)
208#pragma clang diagnostic pop
209#endif
210
211
212
213#endif // TDLS_SOLVERS_TILE_OPERATIONS_HPP
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
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:68
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:137
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:103
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU 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:177
Detection of the compiler, of the language and of the side of the compilation pass: the one place whe...