Compile-time knobs of the ArrowheadLUpp solvers, carrying the tuned defaults. More...
#include <config.hpp>
Public Attributes | |
| ArrowheadCoupling | coupling = ArrowheadCoupling::Dense |
| ArrowheadStorage | storage = ArrowheadStorage::Full |
| MatrixLayout | layout = MatrixLayout::RowMajor |
| StructuralReal< T > | condensation_pivot_threshold = T(0.01) |
| ArrowheadCondensation | condensation = ArrowheadCondensation::MinFlops |
| int | block_tile_size = 4 |
| RowInterchange | block_row_interchange = RowInterchange::Auto |
| StructuralReal< T > | block_relative_pivot_threshold = T(0.01) |
| StructuralReal< T > | block_oot_pivot_threshold = std::is_same_v<T, float> ? T(1e-6f) : T(1e-12) |
| int | border_tile_size = 3 |
| Tile size of TiledLUpp on the border system. | |
| int | rhs_pass_width = -1 |
| bool | gather_rhs = false |
| bool | schur_in_place = false |
| bool | unroll_loops = true |
| bool | streamed_out_of_line = false |
Detailed Description
struct tdls::ArrowheadLUppConfig< T >
Compile-time knobs of the ArrowheadLUpp solvers, carrying the tuned defaults.
- Template Parameters
-
T scalar type (float, double or long double)
Member Data Documentation
◆ block_oot_pivot_threshold
| StructuralReal<T> tdls::ArrowheadLUppConfig< T >::block_oot_pivot_threshold = std::is_same_v<T, float> ? T(1e-6f) : T(1e-12) |
Out-of-tile trigger of TiledLUpp on the blocks (TiledLUppConfig::oot_pivot_threshold), two orders below the default of TiledLUpp: the out-of-tile recovery, a divergent path, fires only on a nearly zero in-tile column.
◆ block_relative_pivot_threshold
| StructuralReal<T> tdls::ArrowheadLUppConfig< T >::block_relative_pivot_threshold = T(0.01) |
Relative threshold of the in-tile pivot choice of TiledLUpp on the blocks (TiledLUppConfig::relative_pivot_threshold), in (0, 1]. The row in place keeps the pivot down to this fraction of the best in-tile candidate: on the blocks of a polycrystal no row moves, so the pivot arrays stay alike across a warp and the reads of the block copy at pivot-dependent offsets stay coalesced. The pivot check of the condensation bounds the multipliers by 1 / condensation_pivot_threshold anyway.
◆ block_row_interchange
| RowInterchange tdls::ArrowheadLUppConfig< T >::block_row_interchange = RowInterchange::Auto |
Row interchanges of TiledLUpp on the diagonal blocks (TiledLUppConfig::row_interchange). Physical eliminates each block in registers: the whole block one tile, the rows interchanged as in LAPACK, every index known at compile time once the loops are unrolled, the right-hand sides [A_kb | y_k] solved the same way; the pivots of one tile of the size of the block (block_relative_pivot_threshold, block_oot_pivot_threshold; block_tile_size and gather_rhs ignored). Logical eliminates it in tiles of block_tile_size, the block on the stack of the thread on GPU, only the working tiles in registers. Auto, the default, takes Physical when the block is small enough (detail::physical_rows_fit): up to blocks of 14 in float and 10 in double on GPU, 8 in float and 12 in double on CPU, long double as double. The pivot check, the recovery data and the second pass are the same either way, and the runtime solver makes the same choice and runs the same operations, so that both solvers stay bitwise identical. Measured gains of the physical interchanges over tiles of 3, blocks that keep their rows, border of 12: 2 to 6 times on the recomputed path, 1.1 to 2.9 times streamed, 1 to 1.6 times stored on GPU (RTX 5080, float and the register budget of double), 1.5 times (median) on CPU up to these sizes. Beyond them the block no longer fits in the registers of a thread and the compilers spill it whole: 5 to 14 times slower from blocks of 16 in float on GPU; on CPU, the tiled solver vectorizes better in float.
◆ block_tile_size
| int tdls::ArrowheadLUppConfig< T >::block_tile_size = 4 |
Tile size of TiledLUpp on the diagonal blocks, when they are not eliminated in registers (block_row_interchange). Tiles of 4 beat tiles of 3 by 1.1 to 1.3 times from blocks of 14 (RTX 5080, float and the register budget of double, the three paths). TiledLUpp searches the pivot inside the tile first: a tile of the size of the block pivots over the whole block, as the physical interchanges do. On blocks that pivot, this refuses far fewer blocks (random blocks of 16, dense coupling: 0.16 refused per system with tiles of 3, 0.01 with one tile) and runs 1.3 to 3 times faster than tiles of 3 at S = 12 and 16 (RTX 5080, float, the matrix stored).
◆ condensation
| ArrowheadCondensation tdls::ArrowheadLUppConfig< T >::condensation = ArrowheadCondensation::MinFlops |
Condensation of an accepted block, see ArrowheadCondensation; it acts under the row-major layout with the dense coupling only. MinFlops, the default, solves the forward stage alone and stores U in D_k for the recovery. MinTraffic runs the whole solve: the backward stage on B more columns, about S^2 B / 2 more multiply-adds per block, but U is neither written by the condensation nor read by the recovery. Measured with MinTraffic against MinFlops on an RTX 5080, in float and with the register budget of double: streamed, 1.46 to 2.2 times faster with blocks of 6, 1.06 to 1.5 times with blocks of 12, the same speed with blocks of 24; stored, 1.1 to 1.6 times faster with the systems one after the other (AoS), but 0.72 to 0.77 times in float with the systems interleaved (SoA) and blocks of 12. In double, streamed: 1.37 to 1.73 times in AoS, 1.53 times in SoA with blocks of 12, but 0.83 times in SoA with blocks of 6. On CPU no trend (0.5 to 1.5 times, slower on small blocks in float). The pivot check is the same, so both refuse the same blocks; the solutions differ by rounding errors.
◆ condensation_pivot_threshold
| StructuralReal<T> tdls::ArrowheadLUppConfig< T >::condensation_pivot_threshold = T(0.01) |
Threshold of the pivot check of the condensation, in [0, 1]. A block is eliminated before the border, its pivots chosen among its own rows only. Each pivot must reach this fraction of the largest candidate of its column, the border rows included, so that no multiplier exceeds 1 / condensation_pivot_threshold in magnitude. A block that fails is refused (a delayed pivot, as in multifrontal solvers) and eliminated in the second pass, where the border rows may serve as pivots. 0 refuses only the singular blocks and those whose multipliers are not finite. The second pass prefers a row of the block as pivot when it reaches this fraction of the largest candidate.
On the jacobians of an implicit Berveiller-Zaoui polycrystal, 0.01 refuses no block. 0.1 bounds the multipliers by 10 instead of 100, but refuses 0.075 % of the blocks under the column-major layout with a border of 12: on GPU, almost every warp then holds a system that takes the second pass (+16 % at 240 phases, RTX 5080). 0.3 refuses 13 to 42 % of them.
◆ coupling
| ArrowheadCoupling tdls::ArrowheadLUppConfig< T >::coupling = ArrowheadCoupling::Dense |
Pattern of the couplings A_bk, see ArrowheadCoupling. The dense coupling accepts every arrowhead matrix; the diagonal one reads and updates block border rows only.
◆ gather_rhs
| bool tdls::ArrowheadLUppConfig< T >::gather_rhs = false |
Permutes the right-hand sides of a block by its pivots through a predicated gather into a second buffer (true) or by cycles in place (false); used where the condensation runs the whole solve (row-major with the diagonal coupling or with MinTraffic, column-major in several passes). It changes the register allocation more than the cost of the permutation: on the RTX 5080, -15 to -20 % with a border of 12 and passes of 2, but up to +18 % with a border of 6 at 240 phases, whether the pivots move or not. To measure per shape.
◆ layout
| MatrixLayout tdls::ArrowheadLUppConfig< T >::layout = MatrixLayout::RowMajor |
Order of the stored matrix: by rows or by columns, for every storage (ColMajor turns Csr into CSC and Bsr into BSC, with blocks stored by columns). The diagonal blocks are factored by TiledLUpp in the same layout, as stored, and the condensation takes the path that suits it (see solver_static.hpp).
◆ rhs_pass_width
| int tdls::ArrowheadLUppConfig< T >::rhs_pass_width = -1 |
Right-hand sides [A_kb | y_k] of a block solved per pass, in [-1, border + 1] for the compile-time solver (the runtime solver takes any width from -1, a width above border + nrhs meaning one pass); 0 solves all border + 1 at once, -1 (the default) chooses: passes of 2 on GPU when the blocks are eliminated in registers, one pass otherwise. With the blocks in registers on GPU, passes of 2 keep fewer values alive: 1.2 to 1.9 times faster than one pass under the row-major layout (median, RTX 5080, float and the register budget of double); under the column-major layout, from 0.8 times (blocks of 10, recomputed) to 15 times (blocks of 14, streamed, where one pass spills the block). On CPU one pass is the best (passes of 2: 0.8 times its speed). With the blocks eliminated by tiles, under the column-major layout, passes gain nothing: one pass reads the pivot check from the forward stage of the solve, several need U and a separate check. The two order the operations differently once a block spans several tiles: a block whose largest multiplier lies within rounding of 1 / condensation_pivot_threshold can pass the check in one pass and fail it in several, and the solutions then differ by rounding errors. A width of 2 or more can also give one pass to an entry point and several to another, whose right-hand sides add columns; with 0, every entry point runs one pass and the same check.
◆ schur_in_place
| bool tdls::ArrowheadLUppConfig< T >::schur_in_place = false |
Accumulates the Schur complement in place, in the border entries of A and y (true), or in a local array of the coupled border rows, alive across the sweep over the blocks (false). In place frees registers but reads and writes these rows for every block: on the RTX 5080 it costs 3 to 28 % under the diagonal coupling and saves 2 to 15 % under the dense one, where the array holds the whole border; neutral on CPU.
◆ storage
| ArrowheadStorage tdls::ArrowheadLUppConfig< T >::storage = ArrowheadStorage::Full |
Storage of the matrix, see ArrowheadStorage. Packed and Csr hold the pattern entries alone, as many in both: the least memory. Bsr holds as many under the dense coupling, whole coupling blocks under the diagonal one.
◆ streamed_out_of_line
| bool tdls::ArrowheadLUppConfig< T >::streamed_out_of_line = false |
Streamed entry points of the compile-time solver kept out of line on GPU (CUDA, HIP, SYCL device code; ignored elsewhere and by the runtime solver). The provider is compiled inside the loop over the blocks, and the GPU inliner may merge the whole solve into a large caller, whose registers it then allocates worse than the two functions apart. Off, the generated code is that of a solver without the knob; on, a polycrystal of 240 grains ran 19 % faster on an RTX 5080. To measure per provider.
◆ unroll_loops
| bool tdls::ArrowheadLUppConfig< T >::unroll_loops = true |
Unroll policy, as in the other families: true, the loops over the entries of a block and of the border carry a forced-unroll pragma, the guard that keeps them in registers on GPU; false, no pragma, for faster compiles. Passed on to the TiledLUpp solvers of the blocks and of the border. Ignored by the runtime solver.
The documentation for this struct was generated from the following file:
- include/tdls/solvers/arrowhead_lupp/config.hpp
Generated by
1.12.0