TDLS 0.4.0
Tiny Device-callable Linear Solvers
Loading...
Searching...
No Matches
storage_map.hpp
Go to the documentation of this file.
1#ifndef TDLS_SOLVERS_ARROWHEAD_LUPP_STORAGE_MAP_HPP
2#define TDLS_SOLVERS_ARROWHEAD_LUPP_STORAGE_MAP_HPP
3
4
5
38
39
40
41#include <tdls/core/macros.hpp>
44
45
46
47namespace tdls {
48
49
50
55template<ArrowheadStorage Storage, ArrowheadCoupling Coupling, MatrixLayout Layout>
58 static constexpr bool diag = Coupling == ArrowheadCoupling::Diagonal;
60 static constexpr bool col = Layout == MatrixLayout::ColMajor;
61
69 [[nodiscard]] TDLS_HOST_DEVICE static constexpr bool valid(const ArrowheadShape& s) noexcept {
70 if (s.border < 1 || s.block < 1 || s.blocks < 1 || s.coupling_row < 0) return false;
71 if (diag ? s.coupling_row > s.border - s.block : s.coupling_row != 0) return false;
72 if (Storage == ArrowheadStorage::Bsr &&
73 (s.border % s.block != 0 || (diag && s.coupling_row % s.block != 0)))
74 return false;
75 if (static_cast<long long>(s.border) + static_cast<long long>(s.block) * s.blocks >=
76 (1LL << 31))
77 return false;
78 return wide_size(s) < (1LL << 31);
79 }
80
86 [[nodiscard]] TDLS_HOST_DEVICE static constexpr long long
87 wide_size(const ArrowheadShape& s) noexcept {
88 const long long B = s.border, S = s.block, NB = s.blocks, n = B + S * NB;
89 if constexpr (Storage == ArrowheadStorage::Full) {
90 return n * n;
91 } else if constexpr (Storage == ArrowheadStorage::Packed) {
92 return B * B + (diag ? NB * S : NB * B * S) + NB * S * B + NB * S * S;
93 } else if constexpr (Storage == ArrowheadStorage::Csr) {
94 if constexpr (col)
95 return B * n + NB * S * (diag ? 1 + S : B + S);
96 else
97 return (diag ? B * B + NB * S : B * n) + NB * S * (B + S);
98 } else {
99 const long long BB = S > 0 ? B / S : 0;
100 if constexpr (col)
101 return (BB * (BB + NB) + NB * (diag ? 2 : BB + 1)) * S * S;
102 else
103 return ((diag ? BB * BB + NB : BB * (BB + NB)) + NB * (BB + 1)) * S * S;
104 }
105 }
106
110 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr int
111 size(const ArrowheadShape& s) noexcept {
112 return static_cast<int>(wide_size(s));
113 }
114
120 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr int
121 abb(const ArrowheadShape& s, const int r, const int c) noexcept {
122 const int B = s.border, S = s.block, NB = s.blocks, n = B + S * NB;
123 if constexpr (Storage == ArrowheadStorage::Full) {
124 return col ? c * n + r : r * n + c;
125 } else if constexpr (Storage == ArrowheadStorage::Packed) {
126 return col ? c * B + r : r * B + c;
127 } else if constexpr (Storage == ArrowheadStorage::Csr) {
128 return col ? c * n + r : csr_row(s, r) + c;
129 } else if constexpr (col) {
130 const int BB = B / S;
131 return ((c / S) * (BB + NB) + r / S) * S * S + in_block(s, r % S, c % S);
132 } else {
133 return (bsr_row(s, r / S) + c / S) * S * S + in_block(s, r % S, c % S);
134 }
135 }
136
144 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr int
145 abk(const ArrowheadShape& s, const int k, const int r, const int i) noexcept {
146 const int B = s.border, S = s.block, NB = s.blocks, n = B + S * NB;
147 if constexpr (Storage == ArrowheadStorage::Full) {
148 return col ? (B + k * S + i) * n + r : r * n + B + k * S + i;
149 } else if constexpr (Storage == ArrowheadStorage::Packed) {
150 return col ? B * B + (k * S + i) * B + r : B * B + (k * B + r) * S + i;
151 } else if constexpr (Storage == ArrowheadStorage::Csr) {
152 return col ? csc_block_col(s, k, i) + r : csr_row(s, r) + B + k * S + i;
153 } else if constexpr (col) {
154 return (bsc_block_col(s, k) + r / S) * S * S + in_block(s, r % S, i);
155 } else {
156 return (bsr_row(s, r / S) + B / S + k) * S * S + in_block(s, r % S, i);
157 }
158 }
159
166 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr int
167 abk_diag(const ArrowheadShape& s, const int k, const int i) noexcept {
168 const int B = s.border, S = s.block, NB = s.blocks, n = B + S * NB, b0 = s.coupling_row;
169 if constexpr (Storage == ArrowheadStorage::Full) {
170 return col ? (B + k * S + i) * n + b0 + i : (b0 + i) * n + B + k * S + i;
171 } else if constexpr (Storage == ArrowheadStorage::Packed) {
172 return B * B + k * S + i;
173 } else if constexpr (Storage == ArrowheadStorage::Csr) {
174 return col ? csc_block_col(s, k, i) : csr_row(s, b0 + i) + B + k;
175 } else if constexpr (col) {
176 return bsc_block_col(s, k) * S * S + in_block(s, i, i);
177 } else {
178 return (bsr_row(s, b0 / S) + B / S + k) * S * S + in_block(s, i, i);
179 }
180 }
181
188 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr int
189 akb(const ArrowheadShape& s, const int k, const int i, const int c) noexcept {
190 const int B = s.border, S = s.block, NB = s.blocks, n = B + S * NB;
191 if constexpr (Storage == ArrowheadStorage::Full) {
192 return col ? c * n + B + k * S + i : (B + k * S + i) * n + c;
193 } else if constexpr (Storage == ArrowheadStorage::Packed) {
194 const int o = B * B + (diag ? NB * S : NB * B * S);
195 return col ? o + (k * B + c) * S + i : o + (k * S + i) * B + c;
196 } else if constexpr (Storage == ArrowheadStorage::Csr) {
197 return col ? c * n + B + k * S + i : csr_block_row(s, k, i) + c;
198 } else if constexpr (col) {
199 const int BB = B / S;
200 return ((c / S) * (BB + NB) + BB + k) * S * S + in_block(s, i, c % S);
201 } else {
202 return (bsr_block_row(s, k) + c / S) * S * S + in_block(s, i, c % S);
203 }
204 }
205
212 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr int
213 d(const ArrowheadShape& s, const int k, const int i, const int j) noexcept {
214 const int B = s.border, S = s.block, NB = s.blocks, n = B + S * NB;
215 if constexpr (Storage == ArrowheadStorage::Full) {
216 return col ? (B + k * S + j) * n + B + k * S + i : (B + k * S + i) * n + B + k * S + j;
217 } else if constexpr (Storage == ArrowheadStorage::Packed) {
218 const int o = B * B + (diag ? NB * S : NB * B * S) + NB * S * B;
219 return o + k * S * S + in_block(s, i, j);
220 } else if constexpr (Storage == ArrowheadStorage::Csr) {
221 return col ? csc_block_col(s, k, j) + (diag ? 1 : B) + i
222 : csr_block_row(s, k, i) + B + j;
223 } else if constexpr (col) {
224 return (bsc_block_col(s, k) + (diag ? 1 : B / S)) * S * S + in_block(s, i, j);
225 } else {
226 return (bsr_block_row(s, k) + B / S) * S * S + in_block(s, i, j);
227 }
228 }
229
236 [[nodiscard]] TDLS_HOST_DEVICE static constexpr bool
237 in_pattern(const ArrowheadShape& s, const int r, const int c) noexcept {
238 const int B = s.border, S = s.block, b0 = s.coupling_row;
239 if (r < B && c < B) return true;
240 if (r < B) return !diag || (r >= b0 && r < b0 + S && (c - B) % S == r - b0);
241 if (c < B) return true;
242 return (r - B) / S == (c - B) / S;
243 }
244
252 TDLS_HOST_DEVICE static constexpr void csr_pattern(const ArrowheadShape& s, int* ptr,
253 int* idx) noexcept {
254 const int n = s.n();
255 int e = 0;
256 for (int a = 0; a < n; ++a) {
257 ptr[a] = e;
258 for (int b = 0; b < n; ++b)
259 if (col ? in_pattern(s, b, a) : in_pattern(s, a, b)) idx[e++] = b;
260 }
261 ptr[n] = e;
262 }
263
272 TDLS_HOST_DEVICE static constexpr void bsr_pattern(const ArrowheadShape& s, int* ptr,
273 int* idx) noexcept {
274 const int BB = s.border / s.block, NB = s.blocks, Rc = s.coupling_row / s.block;
275 auto in = [&](const int R, const int Cb) {
276 if (R < BB) return Cb < BB || !diag || R == Rc;
277 return Cb < BB || Cb == R;
278 };
279 int e = 0;
280 for (int a = 0; a < BB + NB; ++a) {
281 ptr[a] = e;
282 for (int b = 0; b < BB + NB; ++b)
283 if (col ? in(b, a) : in(a, b)) idx[e++] = b;
284 }
285 ptr[BB + NB] = e;
286 }
287
288 private:
291 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr int
292 in_block(const ArrowheadShape& s, const int a, const int b) noexcept {
293 return col ? b * s.block + a : a * s.block + b;
294 }
296 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr int
297 csr_row(const ArrowheadShape& s, const int r) noexcept {
298 const int B = s.border, S = s.block, NB = s.blocks, b0 = s.coupling_row;
299 if constexpr (diag) {
300 const int coupled = r < b0 ? 0 : (r < b0 + S ? r - b0 : S);
301 return r * B + NB * coupled;
302 } else {
303 return r * (B + S * NB);
304 }
305 }
307 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr int
308 csr_block_row(const ArrowheadShape& s, const int k, const int i) noexcept {
309 const int B = s.border, S = s.block, NB = s.blocks;
310 return (diag ? B * B + NB * S : B * (B + S * NB)) + (k * S + i) * (B + S);
311 }
313 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr int
314 csc_block_col(const ArrowheadShape& s, const int k, const int j) noexcept {
315 const int B = s.border, S = s.block, NB = s.blocks;
316 return B * (B + S * NB) + (k * S + j) * (diag ? 1 + S : B + S);
317 }
319 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr int
320 bsr_row(const ArrowheadShape& s, const int R) noexcept {
321 const int BB = s.border / s.block, NB = s.blocks, Rc = s.coupling_row / s.block;
322 if constexpr (diag)
323 return R * BB + (R > Rc ? NB : 0);
324 else
325 return R * (BB + NB);
326 }
328 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr int
329 bsr_block_row(const ArrowheadShape& s, const int k) noexcept {
330 const int BB = s.border / s.block, NB = s.blocks;
331 return (diag ? BB * BB + NB : BB * (BB + NB)) + k * (BB + 1);
332 }
334 [[nodiscard]] TDLS_HOST_DEVICE TDLS_FORCE_INLINE_GPU static constexpr int
335 bsc_block_col(const ArrowheadShape& s, const int k) noexcept {
336 const int BB = s.border / s.block, NB = s.blocks;
337 return BB * (BB + NB) + k * (diag ? 2 : BB + 1);
338 }
339};
340
341
342
343} // namespace tdls
344
345
346
347#endif // TDLS_SOLVERS_ARROWHEAD_LUPP_STORAGE_MAP_HPP
Shape and compile-time configuration of the ArrowheadLUpp solver family.
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_HOST_DEVICE
__host__ __device__ under CUDA, the equivalent attributes under HIP, empty elsewhere.
Definition macros.hpp:50
Shared vocabulary and conventions of the solver families.
Shape of an arrowhead matrix: a template argument of the compile-time solver, a function argument of ...
Definition config.hpp:111
Index of every entry of an arrowhead matrix in its storage.
Definition storage_map.hpp:56
static TDLS_HOST_DEVICE constexpr long long wide_size(const ArrowheadShape &s) noexcept
Number of stored entries, in 64-bit arithmetic (the overflow check of valid()).
Definition storage_map.hpp:87
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int abb(const ArrowheadShape &s, const int r, const int c) noexcept
Border-border entry (r, c), r and c < border.
Definition storage_map.hpp:121
static TDLS_HOST_DEVICE constexpr bool in_pattern(const ArrowheadShape &s, const int r, const int c) noexcept
Whether entry (r, c) of the full matrix belongs to the pattern.
Definition storage_map.hpp:237
static constexpr bool col
Whether the matrix is stored by columns.
Definition storage_map.hpp:60
static TDLS_HOST_DEVICE constexpr void csr_pattern(const ArrowheadShape &s, int *ptr, int *idx) noexcept
Writes the canonical compressed index arrays of the pattern: CSR (row pointers, column indices) under...
Definition storage_map.hpp:252
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int size(const ArrowheadShape &s) noexcept
Number of stored entries.
Definition storage_map.hpp:111
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int abk(const ArrowheadShape &s, const int k, const int r, const int i) noexcept
Coupling entry of border row r and unknown i of block k (dense coupling).
Definition storage_map.hpp:145
static constexpr bool diag
Whether the couplings are diagonal.
Definition storage_map.hpp:58
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int akb(const ArrowheadShape &s, const int k, const int i, const int c) noexcept
Entry of row i of block k in border column c.
Definition storage_map.hpp:189
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int d(const ArrowheadShape &s, const int k, const int i, const int j) noexcept
Entry (i, j) of the diagonal block k.
Definition storage_map.hpp:213
static TDLS_HOST_DEVICE constexpr void bsr_pattern(const ArrowheadShape &s, int *ptr, int *idx) noexcept
Writes the canonical block compressed index arrays, in square blocks of block entries: BSR by block r...
Definition storage_map.hpp:272
static TDLS_HOST_DEVICE constexpr bool valid(const ArrowheadShape &s) noexcept
Whether a shape suits this storage and coupling: positive sizes, coupled rows inside the border (diag...
Definition storage_map.hpp:69
TDLS_HOST_DEVICE static TDLS_FORCE_INLINE_GPU constexpr int abk_diag(const ArrowheadShape &s, const int k, const int i) noexcept
Coupling entry of border row coupling_row + i and unknown i of block k (diagonal coupling).
Definition storage_map.hpp:167