config.hpp

Go to the documentation of this file.

#ifndef TDLS_SOLVERS_ARROWHEAD_LUPP_CONFIG_HPP
#define TDLS_SOLVERS_ARROWHEAD_LUPP_CONFIG_HPP
 
 
 
 
 
 
#include <type_traits>
 
#include <tdls/core/macros.hpp>
#include <tdls/core/structural_real.hpp>
#include <tdls/solvers/options.hpp>
#include <tdls/solvers/tiled_lupp/config.hpp>
 
 
 
namespace tdls {
 
 
 
enum class ArrowheadStorage {
    Full,   
    Packed, 
    Csr,    
    Bsr     
};
 
 
 
enum class ArrowheadCoupling {
    Dense,   
    Diagonal 
};
 
 
 
enum class ArrowheadCondensation {
    MinFlops,  
    MinTraffic 
};
 
 
 
struct ArrowheadShape {
    int border = 1;
    int block = 1;
    int blocks = 1;
    int coupling_row = 0;
 
    [[nodiscard]] TDLS_HOST_DEVICE constexpr int n() const noexcept {
        return border + block * blocks;
    }
};
 
 
 
template<typename T>
struct ArrowheadLUppConfig {
 
    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;
 
    int rhs_pass_width = -1;
 
    bool gather_rhs = false;
 
    bool schur_in_place = false;
 
    bool unroll_loops = true;
 
    bool streamed_out_of_line = false;
};
 
 
 
namespace detail {
 
template<typename T>
[[nodiscard]] constexpr TiledLUppConfig<T>
arrowhead_block_config(const ArrowheadLUppConfig<T>& config) noexcept {
    TiledLUppConfig<T> c;
    c.tile_size                = config.block_tile_size;
    c.relative_pivot_threshold = config.block_relative_pivot_threshold;
    c.oot_pivot_threshold      = config.block_oot_pivot_threshold;
    c.unroll_loops             = config.unroll_loops;
    c.layout                   = config.layout;
    return c;
}
 
template<typename T>
[[nodiscard]] constexpr TiledLUppConfig<T>
arrowhead_physical_block_config(const ArrowheadLUppConfig<T>& config) noexcept {
    TiledLUppConfig<T> c = arrowhead_block_config(config);
    c.row_interchange    = RowInterchange::Physical;
    return c;
}
 
template<typename T>
[[nodiscard]] TDLS_HOST_DEVICE constexpr bool
arrowhead_register_blocks(const ArrowheadLUppConfig<T>& config, const int S) noexcept {
    return physical_rows<T>(config.block_row_interchange, S);
}
 
template<typename T>
[[nodiscard]] TDLS_HOST_DEVICE constexpr int
arrowhead_rhs_pass_width(const ArrowheadLUppConfig<T>& config,
                         [[maybe_unused]] const bool reg) noexcept {
    if (config.rhs_pass_width >= 0) return config.rhs_pass_width;
#if TDLS_TARGET_GPU
    return reg ? 2 : 0;
#else
    return 0;
#endif
}
 
template<typename T>
[[nodiscard]] TDLS_HOST_DEVICE constexpr bool
arrowhead_register_blocks_here(const RowInterchange interchange, const int S) noexcept {
    return physical_rows_here<T>(interchange, S);
}
 
[[nodiscard]] TDLS_HOST_DEVICE constexpr int
arrowhead_rhs_pass_width_here(const int width, [[maybe_unused]] const bool reg) noexcept {
    if (width >= 0) return width;
#if defined(TDLS_DETAIL_SIDE_DISPATCH)
    if (!std::is_constant_evaluated()) return on_device_side() && reg ? 2 : 0;
#endif
#if TDLS_TARGET_GPU
    return reg ? 2 : 0;
#else
    return 0;
#endif
}
 
template<typename T, bool gpu>
[[nodiscard]] constexpr ArrowheadLUppConfig<T>
arrowhead_resolved_config(ArrowheadLUppConfig<T> config, const int S) noexcept {
    const bool reg               = physical_rows_on<T, gpu>(config.block_row_interchange, S);
    config.block_row_interchange = reg ? RowInterchange::Physical : RowInterchange::Logical;
    if (config.rhs_pass_width < 0) config.rhs_pass_width = gpu && reg ? 2 : 0;
    return config;
}
 
template<typename T>
[[nodiscard]] constexpr TiledLUppConfig<T>
arrowhead_border_config(const ArrowheadLUppConfig<T>& config) noexcept {
    TiledLUppConfig<T> c;
    c.tile_size    = config.border_tile_size;
    c.unroll_loops = config.unroll_loops;
    return c;
}
 
} // namespace detail
 
 
 
} // namespace tdls
 
 
 
#endif // TDLS_SOLVERS_ARROWHEAD_LUPP_CONFIG_HPP