Implementations of Kernels and Functions for Matrix Operations on the GPU. More...

#include "../../include/basic_types/codi_host_only.hpp"
#include <algorithm>
#include <cstring>
#include <type_traits>
#include "../../include/linear_algebra/CMatrixInverse.hpp"
#include "../../include/linear_algebra/CSysMatrix.inl"
#include "../../include/linear_algebra/GPUComms.cuh"

Macros
#define	INSTANTIATE_MATRIX(TYPE)

Functions
template<class ScalarType , class QuantType , class QuantScaleType , bool Quantized>
__global__ void	LU_SGS_ForwardKernel (const su2uint __restrict__ level_idx, unsigned long level_begin, unsigned long level_size, unsigned long nVar, DeviceLDU< ScalarType > M, const QuantType __restrict__ q_l, const QuantScaleType __restrict__ q_scale_l, const ScalarType __restrict__ invD, const ScalarType __restrict__ vec, ScalarType __restrict__ prod)
	Exact forward substitution for the rows of one level, x* = D^{-1}.(b-Lx*)

template<class ScalarType , class QuantType , class QuantScaleType , bool Quantized>
__global__ void	LU_SGS_BackwardKernel (const su2uint __restrict__ level_idx, unsigned long level_begin, unsigned long level_size, unsigned long nRows, unsigned long nVar, DeviceLDU< ScalarType > M, const QuantType __restrict__ q_u, const QuantScaleType __restrict__ q_scale_u, const ScalarType __restrict__ invD, ScalarType *__restrict__ prod)
	Exact backward substitution for the rows of one level, x* = D^{-1}.(D.x* - U.x) = x* - D^{-1}.U.x.

Detailed Description

Implementations of Kernels and Functions for Matrix Operations on the GPU.

Author: A. Raj, Jesse Li, P. Gomes

Version: 8.5.0 "Harrier"

SU2 Project Website: https://su2code.github.io

The SU2 Project is maintained by the SU2 Foundation (http://su2foundation.org)

SU2 is free software; you can redistribute it and/or modify it under the terms of the GNU Lesser General Public License as published by the Free Software Foundation; either version 2.1 of the License, or (at your option) any later version.

SU2 is distributed in the hope that it will be useful, but WITHOUT ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU Lesser General Public License for more details.

You should have received a copy of the GNU Lesser General Public License along with SU2. If not, see http://www.gnu.org/licenses/.

Macro Definition Documentation

◆ INSTANTIATE_MATRIX

#define INSTANTIATE_MATRIX ( TYPE )

Value:

template void CSysMatrix<TYPE>::HtDTransfer(bool trigger) const;                            \
template void CSysMatrix<TYPE>::MatrixVectorProductGPU(const CSysVector<TYPE>& vec,         \
                                                       CSysVector<TYPE>& prod,              \
                                                       CGeometry* geometry,                 \
                                                       const CConfig* config) const;        \
template void CSysMatrix<TYPE>::QuantizeDiagonalBlocksGPU();                                \
template void CSysMatrix<TYPE>::BuildJacobiPreconditionerGPU();                             \
template void CSysMatrix<TYPE>::BuildILUPreconditionerGPU();                                \
template void CSysMatrix<TYPE>::BuildLU_SGSPreconditionerGPU();                             \
template void CSysMatrix<TYPE>::ComputeILUPreconditionerGPU(const CSysVector<TYPE>& vec,    \
                                                            CSysVector<TYPE>& prod) const;  \
template void CSysMatrix<TYPE>::ComputeJacobiPreconditionerGPU(const CSysVector<TYPE>& vec, \
                                                               CSysVector<TYPE>& prod,      \
                                                               CGeometry* geometry,         \
                                                               const CConfig* config) const;\
template void CSysMatrix<TYPE>::ComputeLU_SGSForwardGPU(const CSysVector<TYPE>& vec,        \
                                                        CSysVector<TYPE>& prod) const;      \
template void CSysMatrix<TYPE>::ComputeLU_SGSBackwardGPU(CSysVector<TYPE>& prod) const;

Function Documentation

◆ LU_SGS_BackwardKernel()

template<class ScalarType , class QuantType , class QuantScaleType , bool Quantized>

__global__ void LU_SGS_BackwardKernel	(	const su2uint *__restrict__	level_idx,
		unsigned long	level_begin,
		unsigned long	level_size,
		unsigned long	nRows,
		unsigned long	nVar,
		DeviceLDU< ScalarType >	M,
		const QuantType *__restrict__	q_u,
		const QuantScaleType *__restrict__	q_scale_u,
		const ScalarType *__restrict__	invD,
		ScalarType *__restrict__	prod
	)

Exact backward substitution for the rows of one level, x* = D^{-1}.(D.x* - U.x) = x* - D^{-1}.U.x.

Note: See notes in IluBackwardKernel for more details

◆ LU_SGS_ForwardKernel()