1#ifndef TDLS_SOLVERS_TILE_OPERATIONS_HPP
2#define TDLS_SOLVERS_TILE_OPERATIONS_HPP
39#if defined(TDLS_DETAIL_COMPILER_CLANG)
40#pragma clang diagnostic push
41#pragma clang diagnostic ignored "-Wpass-failed"
53template<
typename T,
int tile_size,
bool unroll_loops>
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) {
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;
83 for (
int row = 0; row < k_extent; ++row) {
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;
101 template<
int diag_extent,
int col_extent>
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];
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];
135 template<
int diag_extent,
int row_extent>
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;
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;
175 template<
int row_extent,
int col_extent,
int k_extent>
178 if constexpr (unroll_loops) {
180 for (
int i = 0; i < row_extent; ++i) {
182 for (
int j = 0; j < col_extent; ++j) {
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;
191 for (
int i = 0; i < row_extent; ++i) {
192 for (
int j = 0; j < col_extent; ++j) {
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;
207#if defined(TDLS_DETAIL_COMPILER_CLANG)
208#pragma clang diagnostic pop
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...