Opm::gpuistl::detail Namespace Reference

Contains wrappers to make the CuBLAS library behave as a modern C++ library with function overlading. More...

Classes
class	CuBlasHandle
	The CuBlasHandle class provides a singleton for the simulator universal cuBlasHandle. More...
class	CuSparseHandle
	The CuSparseHandle class provides a singleton for the simulator universal cuSparseHandle. More...
class	CuSparseResource
	The CuSparseResource class wraps a CuSparse resource in a proper RAII pattern. More...
class	FlexibleSolverWrapper
	FlexibleSolverWrapper is compilational trick to reduce compile time overhead. More...
class	has_should_call_pre
	The has_should_call_pre class detects the presence of the method shouldCallPre. More...
class	has_should_call_post
	The has_should_call_post class detects the presence of the method shouldCallPost. More...
class	has_communication
	The has_communication class checks if the type has the member function getCommunication. More...
class	is_a_well_operator
	The is_a_well_operator class tries to guess if the operator is a well type operator. More...

Typedefs
using	GpuSparseMatrixDescription = CuSparseResource<cusparseMatDescr_t>
	GpuSparseMatrixDescription holder.
using	GpuSparseMatrixDescriptionPtr = std::shared_ptr<CuSparseResource<cusparseMatDescr_t>>
	Pointer to GpuSparseMatrixDescription holder.

Functions
template<typename func>
int	tuneThreadBlockSize (func &f, std::string descriptionOfFunction)
	Function that tests the best thread block size, assumes the provided function depends on threadblock-size.
std::vector< int >	createReorderedToNatural (const Opm::SparseTable< size_t > &levelSets)
std::vector< int >	createNaturalToReordered (const Opm::SparseTable< size_t > &levelSets)
template<class M, class field_type, class GPUM>
std::unique_ptr< GPUM >	createReorderedMatrix (const M &naturalMatrix, const std::vector< int > &reorderedToNatural)
template<class M, class field_type, class GPUM>
std::tuple< std::unique_ptr< GPUM >, std::unique_ptr< GPUM > >	extractLowerAndUpperMatrices (const M &naturalMatrix, const std::vector< int > &reorderedToNatural)
template<typename T, bool transpose>
void	getQuasiImpesWeights (const GpuSparseMatrixWrapper< T > &matrix, std::size_t pressureVarIndex, GpuVector< T > &weights, const GpuVector< int > &diagonalIndices)
	Calculates quasi-IMPES weights for CPR preconditioner on GPU.
template<typename T, bool transpose>
void	calculateCoarseEntries (const GpuSparseMatrixWrapper< T > &fineMatrix, GpuSparseMatrixWrapper< T > &coarseMatrix, const GpuVector< T > &weights, std::size_t pressureVarIndex)
	Calculates the coarse level matrix entries based on the fine level matrix and weights.
template<typename T, bool transpose>
void	restrictVector (const GpuVector< T > &fine, GpuVector< T > &coarse, const GpuVector< T > &weights, std::size_t pressureVarIndex)
	Restricts a fine level vector to a coarse level vector based on pressure index.
template<typename T, bool transpose>
void	prolongateVector (const GpuVector< T > &coarse, GpuVector< T > &fine, const GpuVector< T > &weights, std::size_t pressureVarIndex)
	Prolongs a coarse level vector to a fine level vector based on pressure index.
std::string	getCublasErrorMessage (cublasStatus_t error, const std::string_view &expression, const std::string_view &filename, const std::string_view &functionName, size_t lineNumber)
	getCublasErrorMessage generates the error message to display for a given error.
void	cublasSafeCall (cublasStatus_t error, const std::string_view &expression, const std::string_view &filename, const std::string_view &functionName, size_t lineNumber)
	cublasSafeCall checks the return type of the CUBLAS expression (function call) and throws an exception if it does not equal CUBLAS_STATUS_SUCCESS.
cublasStatus_t	cublasWarnIfError (cublasStatus_t error, const std::string_view &expression, const std::string_view &filename, const std::string_view &functionName, size_t lineNumber)
	cublasWarnIfError checks the return type of the CUBLAS expression (function call) and issues a warning if it does not equal CUBLAS_STATUS_SUCCESS.
cublasStatus_t	cublasScal (cublasHandle_t handle, int n, const double *alpha, double *x, int incx)
cublasStatus_t	cublasScal (cublasHandle_t handle, int n, const float *alpha, float *x, int incx)
cublasStatus_t	cublasScal (cublasHandle_t handle, int n, const int *alpha, int *x, int incx)
cublasStatus_t	cublasAxpy (cublasHandle_t handle, int n, const double *alpha, const double *x, int incx, double *y, int incy)
cublasStatus_t	cublasAxpy (cublasHandle_t handle, int n, const float *alpha, const float *x, int incx, float *y, int incy)
cublasStatus_t	cublasAxpy (cublasHandle_t handle, int n, const int *alpha, const int *x, int incx, int *y, int incy)
cublasStatus_t	cublasDot (cublasHandle_t handle, int n, const double *x, int incx, const double *y, int incy, double *result)
cublasStatus_t	cublasDot (cublasHandle_t handle, int n, const float *x, int incx, const float *y, int incy, float *result)
cublasStatus_t	cublasDot (cublasHandle_t handle, int n, const int *x, int incx, const int *y, int incy, int *result)
cublasStatus_t	cublasNrm2 (cublasHandle_t handle, int n, const double *x, int incx, double *result)
cublasStatus_t	cublasNrm2 (cublasHandle_t handle, int n, const float *x, int incx, float *result)
cublasStatus_t	cublasNrm2 (cublasHandle_t handle, int n, const int *x, int incx, int *result)
GpuSparseMatrixDescriptionPtr	createMatrixDescription ()
	createMatrixDescription creates a default matrix description
GpuSparseMatrixDescriptionPtr	createLowerDiagonalDescription ()
	createLowerDiagonalDescription creates a lower diagonal matrix description
GpuSparseMatrixDescriptionPtr	createUpperDiagonalDescription ()
	createUpperDiagonalDescription creates an upper diagonal matrix description
std::string	getCusparseErrorCodeToString (int code)
	getCusparseErrorCodeToString Converts an error code returned from a cusparse function a human readable string.
std::string	getCusparseErrorMessage (cusparseStatus_t error, const std::string_view &expression, const std::string_view &filename, const std::string_view &functionName, size_t lineNumber)
	getCusparseErrorMessage generates the error message to display for a given error.
void	cusparseSafeCall (cusparseStatus_t error, const std::string_view &expression, const std::string_view &filename, const std::string_view &functionName, size_t lineNumber)
	cusparseSafeCall checks the return type of the CUSPARSE expression (function call) and throws an exception if it does not equal CUSPARSE_STATUS_SUCCESS.
cusparseStatus_t	cusparseWarnIfError (cusparseStatus_t error, const std::string_view &expression, const std::string_view &filename, const std::string_view &functionName, size_t lineNumber)
	cusparseWarnIfError checks the return type of the CUSPARSE expression (function call) and issues a warning if it does not equal CUSPARSE_STATUS_SUCCESS.
cusparseStatus_t	cusparseBsrilu02_analysis (cusparseHandle_t handle, cusparseDirection_t dirA, int mb, int nnzb, const cusparseMatDescr_t descrA, double *bsrSortedVal, const int *bsrSortedRowPtr, const int *bsrSortedColInd, int blockDim, bsrilu02Info_t info, cusparseSolvePolicy_t policy, void *pBuffer)
cusparseStatus_t	cusparseBsrsv2_analysis (cusparseHandle_t handle, cusparseDirection_t dirA, cusparseOperation_t transA, int mb, int nnzb, const cusparseMatDescr_t descrA, const double *bsrSortedValA, const int *bsrSortedRowPtrA, const int *bsrSortedColIndA, int blockDim, bsrsv2Info_t info, cusparseSolvePolicy_t policy, void *pBuffer)
cusparseStatus_t	cusparseBsrsv2_analysis (cusparseHandle_t handle, cusparseDirection_t dirA, cusparseOperation_t transA, int mb, int nnzb, const cusparseMatDescr_t descrA, const float *bsrSortedValA, const int *bsrSortedRowPtrA, const int *bsrSortedColIndA, int blockDim, bsrsv2Info_t info, cusparseSolvePolicy_t policy, void *pBuffer)
cusparseStatus_t	cusparseBsrilu02_analysis (cusparseHandle_t handle, cusparseDirection_t dirA, int mb, int nnzb, const cusparseMatDescr_t descrA, float *bsrSortedVal, const int *bsrSortedRowPtr, const int *bsrSortedColInd, int blockDim, bsrilu02Info_t info, cusparseSolvePolicy_t policy, void *pBuffer)
cusparseStatus_t	cusparseBsrsv2_solve (cusparseHandle_t handle, cusparseDirection_t dirA, cusparseOperation_t transA, int mb, int nnzb, const double *alpha, const cusparseMatDescr_t descrA, const double *bsrSortedValA, const int *bsrSortedRowPtrA, const int *bsrSortedColIndA, int blockDim, bsrsv2Info_t info, const double *f, double *x, cusparseSolvePolicy_t policy, void *pBuffer)
cusparseStatus_t	cusparseBsrsv2_solve (cusparseHandle_t handle, cusparseDirection_t dirA, cusparseOperation_t transA, int mb, int nnzb, const float *alpha, const cusparseMatDescr_t descrA, const float *bsrSortedValA, const int *bsrSortedRowPtrA, const int *bsrSortedColIndA, int blockDim, bsrsv2Info_t info, const float *f, float *x, cusparseSolvePolicy_t policy, void *pBuffer)
cusparseStatus_t	cusparseBsrilu02_bufferSize (cusparseHandle_t handle, cusparseDirection_t dirA, int mb, int nnzb, const cusparseMatDescr_t descrA, double *bsrSortedVal, const int *bsrSortedRowPtr, const int *bsrSortedColInd, int blockDim, bsrilu02Info_t info, int *pBufferSizeInBytes)
cusparseStatus_t	cusparseBsrilu02_bufferSize (cusparseHandle_t handle, cusparseDirection_t dirA, int mb, int nnzb, const cusparseMatDescr_t descrA, float *bsrSortedVal, const int *bsrSortedRowPtr, const int *bsrSortedColInd, int blockDim, bsrilu02Info_t info, int *pBufferSizeInBytes)
cusparseStatus_t	cusparseBsrsv2_bufferSize (cusparseHandle_t handle, cusparseDirection_t dirA, cusparseOperation_t transA, int mb, int nnzb, const cusparseMatDescr_t descrA, double *bsrSortedValA, const int *bsrSortedRowPtrA, const int *bsrSortedColIndA, int blockDim, bsrsv2Info_t info, int *pBufferSizeInBytes)
cusparseStatus_t	cusparseBsrsv2_bufferSize (cusparseHandle_t handle, cusparseDirection_t dirA, cusparseOperation_t transA, int mb, int nnzb, const cusparseMatDescr_t descrA, float *bsrSortedValA, const int *bsrSortedRowPtrA, const int *bsrSortedColIndA, int blockDim, bsrsv2Info_t info, int *pBufferSizeInBytes)
cusparseStatus_t	cusparseBsrilu02 (cusparseHandle_t handle, cusparseDirection_t dirA, int mb, int nnzb, const cusparseMatDescr_t descrA, double *bsrSortedVal, const int *bsrSortedRowPtr, const int *bsrSortedColInd, int blockDim, bsrilu02Info_t info, cusparseSolvePolicy_t policy, void *pBuffer)
cusparseStatus_t	cusparseBsrilu02 (cusparseHandle_t handle, cusparseDirection_t dirA, int mb, int nnzb, const cusparseMatDescr_t descrA, float *bsrSortedVal, const int *bsrSortedRowPtr, const int *bsrSortedColInd, int blockDim, bsrilu02Info_t info, cusparseSolvePolicy_t policy, void *pBuffer)
cusparseStatus_t	cusparseBsrmv (cusparseHandle_t handle, cusparseDirection_t dirA, cusparseOperation_t transA, int mb, int nb, int nnzb, const double *alpha, const cusparseMatDescr_t descrA, const double *bsrSortedValA, const int *bsrSortedRowPtrA, const int *bsrSortedColIndA, int blockDim, const double *x, const double *beta, double *y)
cusparseStatus_t	cusparseBsrmv (cusparseHandle_t handle, cusparseDirection_t dirA, cusparseOperation_t transA, int mb, int nb, int nnzb, const float *alpha, const cusparseMatDescr_t descrA, const float *bsrSortedValA, const int *bsrSortedRowPtrA, const int *bsrSortedColIndA, int blockDim, const float *x, const float *beta, float *y)
template<class Matrix>
const Matrix	makeMatrixWithNonzeroDiagonal (const Matrix &matrix, const typename Matrix::field_type replacementValue=std::numeric_limits< typename Matrix::field_type >::epsilon())
	makeMatrixWithNonzeroDiagonal creates a new matrix with the zero diagonal elements (when viewed as a matrix of scalrars) set to replacementValue
template<class Operator, class BlockType>
Dune::BCRSMatrix< BlockType >	makeCPUMatrix (const Operator &op)
	Utility functions for GPU preconditioner creation.
std::string	getCudaErrorMessage (cudaError_t error, const std::string_view &expression, const std::string_view &filename, const std::string_view &functionName, size_t lineNumber)
	getCudaErrorMessage generates the error message to display for a given error.
void	cudaSafeCall (cudaError_t error, const std::string_view &expression, const std::string_view &filename, const std::string_view &functionName, size_t lineNumber)
	cudaSafeCall checks the return type of the GPU expression (function call) and throws an exception if it does not equal cudaSuccess.
void	cudaWarnIfError (cudaError_t error, const std::string_view &expression, const std::string_view &filename, const std::string_view &functionName, size_t lineNumber)
	cudaWarnIfError checks the return type of the GPU expression (function call) and issues a warning if it does not equal cudaSuccess.
template<class T, int blocksize>
void	copyMatDataToReordered (const T *srcMatrix, const int *srcRowIndices, T *dstMatrix, int *dstRowIndices, int *naturalToReordered, size_t numberOfRows, int threadBlockSize)
	Reorders the elements of a matrix by copying them from one matrix to another using a permutation list.
template<class T, int blocksize>
void	copyMatDataToReorderedSplit (const T *srcMatrix, const int *srcRowIndices, const int *srcColumnIndices, T *dstLowerMatrix, int *dstLowerRowIndices, T *dstUpperMatrix, int *dstUpperRowIndices, T *dstDiag, int *naturalToReordered, size_t numberOfRows, int threadBlockSize)
	Reorders the elements of a matrix by copying them from one matrix to a split matrix using a permutation list.
auto	makeSafeVectorDescriptor ()
	Create RAII-managed cuSPARSE dense vector descriptor.
auto	makeSafeMatrixDescriptor ()
	Create RAII-managed cuSPARSE sparse matrix descriptor.
template<class T, class M>
std::vector< T >	extractNonzeroValues (const M &matrix)
	Extract non-zero values from a DUNE matrix in the order expected by cuSPARSE.
template<class M>
std::pair< std::vector< int >, std::vector< int > >	extractSparsityPattern (const M &matrix)
	Extract sparsity pattern from a DUNE matrix.
template<class T, class M>
T *	findFirstElementPointer (const M &matrix)
	Find pointer to first matrix element for direct memory access.
template<class M>
void	validateMatrixConversion (const M &matrix, const std::vector< int > &columnIndices, const std::vector< int > &rowIndices)
	Validate matrix conversion parameters.
void	validateMatrixCompatibility (size_t nonzeroes1, size_t blockSize1, size_t N1, size_t nonzeroes2, size_t blockSize2, size_t N2)
	Validate that two matrices have compatible dimensions.
void	validateVectorMatrixSizes (size_t vectorSize, size_t matrixBlockSize, size_t matrixColumns)
	Validate vector-matrix size compatibility for operations.
constexpr size_t	getThreads (size_t numberOfRows)
size_t	getBlocks (size_t numberOfRows)
template<class Kernel>
int	getCudaRecomendedThreadBlockSize (Kernel k, int suggestedThrBlockSize=-1)
int	getNumberOfBlocks (int wantedThreads, int threadBlockSize)
template<class T>
bool	isGPUPointer (const T *ptr)
	Checks whether the given pointer is associated with GPU device memory.
template<class T, class D>
bool	isGPUPointer (const std::unique_ptr< T, D > &ptr)
	Checks if the given std::unique_ptr with custom deleter refers to GPU memory.
template<class T>
bool	isGPUPointer (const std::shared_ptr< T > &ptr)
	Checks if the given std::shared_ptr refers to GPU memory.
bool	isValidMatrixStorageMPScheme (int scheme)
__host__ __device__ constexpr bool	storeDiagonalAsFloat (MatrixStorageMPScheme scheme)
__host__ __device__ constexpr bool	storeOffDiagonalAsFloat (MatrixStorageMPScheme scheme)
__host__ __device__ constexpr bool	usingMixedPrecision (MatrixStorageMPScheme scheme)
template<class PreconditionerType>
constexpr bool	shouldCallPreconditionerPre ()
	Tests (compile time) if the preconditioner type needs to call pre() before a call to apply().
template<class PreconditionerType>
constexpr bool	shouldCallPreconditionerPost ()
	Tests (compile time) if the preconditioner type needs to call post() after a call to apply(...).
template<class T>
void	computeDiagIndicesNoReorder (const int *rowIndices, const int *colIndices, const size_t *indexConversion, int rows, size_t *diagIndices)
	Computes indices of diagonal elements for non-reordered GPU preconditioner.
template<class T>
void	computeDiagIndices (const int *rowIndices, const int *colIndices, const int *reorderedToNatural, int rows, size_t *diagIndices)
	Computes indices of diagonal elements for reordered GPU preconditioner.
int	to_int (std::size_t s)
	to_int converts a (on most relevant platforms) 64 bits unsigned size_t to a signed 32 bits signed int
__host__ __device__ std::size_t	to_size_t (int i)
	to_size_t converts a (on most relevant platforms) a 32 bit signed int to a 64 bits unsigned int
template<class T>
void	setVectorValue (T *deviceData, size_t numberOfElements, const T &value)
	setVectorValue sets every element of deviceData to value
template<class T>
void	setZeroAtIndexSet (T *deviceData, size_t numberOfElements, const int *indices)
	setZeroAtIndexSet sets deviceData to zero in the indices of contained in indices
template<class T>
T	innerProductAtIndices (cublasHandle_t cublasHandle, const T *deviceA, const T *deviceB, T *buffer, size_t numberOfElements, const int *indices)
	innerProductAtIndices computes the inner product between deviceA[indices] and deviceB[indices]
template<class T>
void	prepareSendBuf (const T *deviceA, T *buffer, size_t numberOfElements, const int *indices)
template<class T>
void	syncFromRecvBuf (T *deviceA, T *buffer, size_t numberOfElements, const int *indices)
template<class T>
void	weightedDiagMV (const T *squareBlockVector, const size_t numberOfRows, const size_t blocksize, T relaxationFactor, const T *srcVec, T *dstVec)
	Compue the weighted matrix vector product where the matrix is diagonal, the diagonal is a vector, meaning we compute the Hadamard product.

Variables
constexpr cudaStream_t	DEFAULT_STREAM = 0
	The default GPU stream (stream 0).
constexpr auto	CUSPARSE_MATRIX_ORDER = CUSPARSE_DIRECTION_ROW
	Default matrix ordering for CuSparse operations.

Detailed Description

Contains wrappers to make the CuBLAS library behave as a modern C++ library with function overlading.

Provides various utilities for doing signed to unsigned conversion, unsigned to signed, 32 bits to 64 bits and 64 bits to 32 bits.

Simple utility structs to test for the existence of functions in types.

Contains wrappers to make the CuSPARSE library behave as a modern C++ library with function overlading.

In simple terms, this allows one to call say cublasScal on both double and single precisision, instead of calling cublasDscal and cublasSscal respectively.

In simple terms, this allows one to call say cusparseBsrilu02_analysis on both double and single precisision, instead of calling cusparseDbsrilu02_analysis and cusparseDbsrilu02_analysis respectively.

Note that there are alternatives to this, see for instance https://stackoverflow.com/questions/257288/templated-check-for-the-existence-of-a-class-member-function , however, this is by far the cleanest approach for where this is going to be used for now.

TODO: Use the requires-keyword once C++20 becomes availble (https://en.cppreference.com/w/cpp/language/constraints ). With C++20 this file can be removed.

The main use case within cuistl is that the cusparse library requires signed int for all its size parameters, while Dune::BlockVector (and relatives) use unsigned size_t.

Typedef Documentation

GpuSparseMatrixDescription

using Opm::gpuistl::detail::GpuSparseMatrixDescription = CuSparseResource<cusparseMatDescr_t>

GpuSparseMatrixDescription holder.

This is internal information needed for most calls to the CuSparse API.

GpuSparseMatrixDescriptionPtr

using Opm::gpuistl::detail::GpuSparseMatrixDescriptionPtr = std::shared_ptr<CuSparseResource<cusparseMatDescr_t>>

Pointer to GpuSparseMatrixDescription holder.

This is internal information needed for most calls to the CuSparse API.

Function Documentation

calculateCoarseEntries()

template<typename T, bool transpose>

void Opm::gpuistl::detail::calculateCoarseEntries	(	const GpuSparseMatrixWrapper< T > &	fineMatrix,
		GpuSparseMatrixWrapper< T > &	coarseMatrix,
		const GpuVector< T > &	weights,
		std::size_t	pressureVarIndex )

Calculates the coarse level matrix entries based on the fine level matrix and weights.

Parameters

fineMatrix	The fine level matrix
coarseMatrix	The coarse level matrix to fill
weights	The weights vector for pressure calculation
pressureVarIndex	The index of the pressure variable

Template Parameters

transpose

Whether to use the transpose representation or not

computeDiagIndices()

template<class T>

void Opm::gpuistl::detail::computeDiagIndices	(	const int *	rowIndices,
		const int *	colIndices,
		const int *	reorderedToNatural,
		int	rows,
		size_t *	diagIndices )

Computes indices of diagonal elements for reordered GPU preconditioner.

Template Parameters

T	Matrix element type

Parameters

[in]	rowIndices	Row indices of the BCSR matrix
[in]	colIndices	Column indices of the BCSR matrix
[in]	reorderedToNatural	Table for reordering to natural numbering
[in]	rows	Number of rows in the matrix
[out]	diagIndices	Array to store the indices of diagonal elements

computeDiagIndicesNoReorder()

template<class T>

void Opm::gpuistl::detail::computeDiagIndicesNoReorder	(	const int *	rowIndices,
		const int *	colIndices,
		const size_t *	indexConversion,
		int	rows,
		size_t *	diagIndices )

Computes indices of diagonal elements for non-reordered GPU preconditioner.

Template Parameters

T	Matrix element type

Parameters

[in]	rowIndices	Row indices of the BCSR matrix
[in]	colIndices	Column indices of the BCSR matrix
[in]	indexConversion	Mapping from levelset-row index to matrix row
[in]	rows	Number of rows in the matrix
[out]	diagIndices	Array to store the indices of diagonal elements

copyMatDataToReordered()

template<class T, int blocksize>

void Opm::gpuistl::detail::copyMatDataToReordered	(	const T *	srcMatrix,
		const int *	srcRowIndices,
		T *	dstMatrix,
		int *	dstRowIndices,
		int *	naturalToReordered,
		size_t	numberOfRows,
		int	threadBlockSize )

Reorders the elements of a matrix by copying them from one matrix to another using a permutation list.

Parameters

	srcMatrix	The source matrix we will copy data from
	srcRowIndices	Pointer to vector on GPU containing row indices for the source matrix compliant wiht bsr format
[out]	dstMatrix	The destination matrix that we copy data to
	dstRowIndices	Pointer to vector on GPU containing riw indices for the destination matrix compliant wiht bsr format
	naturalToReordered	Permuation list that converts indices in the src matrix to the indices in the dst matrix
	numberOfRows	The number of rows in the matrices
	threadBlockSize	Block size for threads

copyMatDataToReorderedSplit()

template<class T, int blocksize>

void Opm::gpuistl::detail::copyMatDataToReorderedSplit	(	const T *	srcMatrix,
		const int *	srcRowIndices,
		const int *	srcColumnIndices,
		T *	dstLowerMatrix,
		int *	dstLowerRowIndices,
		T *	dstUpperMatrix,
		int *	dstUpperRowIndices,
		T *	dstDiag,
		int *	naturalToReordered,
		size_t	numberOfRows,
		int	threadBlockSize )

Reorders the elements of a matrix by copying them from one matrix to a split matrix using a permutation list.

Parameters

	srcMatrix	The source matrix we will copy data from
	srcRowIndices	Pointer to vector on GPU containing row indices for the source matrix compliant with bsr format
	srcColumnIndices	Pointer to vector on GPU containing column indices for the source matrix compliant with bsr format
[out]	dstLowerMatrix	The destination of entries that originates from the strictly lower triangular matrix
	dstLowerRowIndices	Pointer to vector on GPU containing rww indices for the destination lower matrix compliant wiht bsr format
[out]	dstUpperMatrix	The destination of entries that originates from the strictly upper triangular matrix
	dstUpperRowIndices	Pointer to vector on GPU containing riw indices for the destination upper matrix compliant wiht bsr format
[out]	dstDiag	The destination buffer for the diagonal part of the matrix
	naturalToReordered	Permuation list that converts indices in the src matrix to the indices in the dst matrix
	numberOfRows	The number of rows in the matrices
	threadBlockSize	Block size for threads

createLowerDiagonalDescription()

GpuSparseMatrixDescriptionPtr Opm::gpuistl::detail::createLowerDiagonalDescription ( )

inline

createLowerDiagonalDescription creates a lower diagonal matrix description

Returns

a lower diagonal matrix description overlapped with options from Opm::gpuistl::detail::createMatrixDescription()

Note

This will assume it has a unit diagonal

createMatrixDescription()

GpuSparseMatrixDescriptionPtr Opm::gpuistl::detail::createMatrixDescription ( )

inline

createMatrixDescription creates a default matrix description

Returns

a matrix description to a general sparse matrix with zero based indexing.

createUpperDiagonalDescription()

GpuSparseMatrixDescriptionPtr Opm::gpuistl::detail::createUpperDiagonalDescription ( )

inline

createUpperDiagonalDescription creates an upper diagonal matrix description

Returns

an upper diagonal matrix description overlapped with options from Opm::gpuistl::detail::createMatrixDescription()

Note

This will assume it has a non-unit diagonal.

cublasSafeCall()

void Opm::gpuistl::detail::cublasSafeCall ( cublasStatus_t error, const std::string_view & expression, const std::string_view & filename, const std::string_view & functionName, size_t lineNumber )

inline

cublasSafeCall checks the return type of the CUBLAS expression (function call) and throws an exception if it does not equal CUBLAS_STATUS_SUCCESS.

Parameters

error	the error code from cublas
expression	the expresison (say "cublasCreate(&handle)")
filename	the code file the error occured in (typically FILE)
functionName	name of the function the error occured in (typically func)
lineNumber	the line number the error occured in (typically LINE)

Example usage:

#include <opm/simulators/linalg/gpuistl/detail/cublas_safe_call.hpp>
#include <cublas_v2.h>
 
void some_function() {
    cublasHandle_t cublasHandle;
    cudaSafeCall(cublasCreate(&cublasHandle), "cublasCreate(&cublasHandle)", __FILE__, __func__, __LINE__);
}

Note

It is probably easier to use the macro OPM_CUBLAS_SAFE_CALL

Todo

Refactor to use std::source_location once we shift to C++20

cublasWarnIfError()

cublasStatus_t Opm::gpuistl::detail::cublasWarnIfError ( cublasStatus_t error, const std::string_view & expression, const std::string_view & filename, const std::string_view & functionName, size_t lineNumber )

inline

cublasWarnIfError checks the return type of the CUBLAS expression (function call) and issues a warning if it does not equal CUBLAS_STATUS_SUCCESS.

Parameters

error	the error code from cublas
expression	the expresison (say "cublasCreate(&handle)")
filename	the code file the error occured in (typically FILE)
functionName	name of the function the error occured in (typically func)
lineNumber	the line number the error occured in (typically LINE)

Returns

the error sent in (for convenience).

Example usage:

#include <opm/simulators/linalg/gpuistl/detail/cublas_safe_call.hpp>
#include <cublas_v2.h>
 
void some_function() {
    cublasHandle_t cublasHandle;
    cublasWarnIfError(cublasCreate(&cublasHandle), "cublasCreate(&cublasHandle)", __FILE__, __func__, __LINE__);
}

Note

It is probably easier to use the macro OPM_CUBLAS_WARN_IF_ERROR

Prefer the cublasSafeCall/OPM_CUBLAS_SAFE_CALL counterpart unless you really don't want to throw an exception.

Todo

Refactor to use std::source_location once we shift to C++20

cudaSafeCall()

void Opm::gpuistl::detail::cudaSafeCall ( cudaError_t error, const std::string_view & expression, const std::string_view & filename, const std::string_view & functionName, size_t lineNumber )

inline

cudaSafeCall checks the return type of the GPU expression (function call) and throws an exception if it does not equal cudaSuccess.

Example usage:

#include <opm/simulators/linalg/gpuistl/detail/gpu_safe_call.hpp>
#include <cuda_runtime.h>
 
void some_function() {
    void* somePointer;
    cudaSafeCall(cudaMalloc(&somePointer, 1), "cudaMalloc(&somePointer, 1)", __FILE__, __func__, __LINE__);
}

Note

It is probably easier to use the macro OPM_GPU_SAFE_CALL

Todo

Refactor to use std::source_location once we shift to C++20

cudaWarnIfError()

void Opm::gpuistl::detail::cudaWarnIfError ( cudaError_t error, const std::string_view & expression, const std::string_view & filename, const std::string_view & functionName, size_t lineNumber )

inline

cudaWarnIfError checks the return type of the GPU expression (function call) and issues a warning if it does not equal cudaSuccess.

Parameters

error	the error code from cublas
expression	the expresison (say "cudaMalloc(&pointer, 1)")
filename	the code file the error occured in (typically FILE)
functionName	name of the function the error occured in (typically func)
lineNumber	the line number the error occured in (typically LINE)

Example usage:

#include <opm/simulators/linalg/gpuistl/detail/gpu_safe_call.hpp>
#include <cuda_runtime.h>
 
void some_function() {
    void* somePointer;
    cudaWarnIfError(cudaMalloc(&somePointer, 1), "cudaMalloc(&somePointer, 1)", __FILE__, __func__, __LINE__);
}

Note

It is probably easier to use the macro OPM_GPU_WARN_IF_ERROR

Prefer the cudaSafeCall/OPM_GPU_SAFE_CALL counterpart unless you really don't want to throw an exception.

Todo

Refactor to use std::source_location once we shift to C++20

cusparseSafeCall()

void Opm::gpuistl::detail::cusparseSafeCall ( cusparseStatus_t error, const std::string_view & expression, const std::string_view & filename, const std::string_view & functionName, size_t lineNumber )

inline

cusparseSafeCall checks the return type of the CUSPARSE expression (function call) and throws an exception if it does not equal CUSPARSE_STATUS_SUCCESS.

Example usage:

#include <opm/simulators/linalg/gpuistl/detail/cusparse_safe_call.hpp>
#include <cublas_v2.h>
 
void some_function() {
    cusparseHandle_t cusparseHandle;
    cusparseSafeCall(cusparseCreate(&cusparseHandle), "cusparseCreate(&cusparseHandle)", __FILE__, __func__,
__LINE__);
}

Note

It is probably easier to use the macro OPM_CUBLAS_SAFE_CALL

Todo

Refactor to use std::source_location once we shift to C++20

cusparseWarnIfError()

cusparseStatus_t Opm::gpuistl::detail::cusparseWarnIfError ( cusparseStatus_t error, const std::string_view & expression, const std::string_view & filename, const std::string_view & functionName, size_t lineNumber )

inline

cusparseWarnIfError checks the return type of the CUSPARSE expression (function call) and issues a warning if it does not equal CUSPARSE_STATUS_SUCCESS.

Parameters

error	the error code from cublas
expression	the expresison (say "cublasCreate(&handle)")
filename	the code file the error occured in (typically FILE)
functionName	name of the function the error occured in (typically func)
lineNumber	the line number the error occured in (typically LINE)

Returns

the error sent in (for convenience).

Example usage:

#include <opm/simulators/linalg/gpuistl/detail/cusparse_safe_call.hpp>
#include <cublas_v2.h>
 
void some_function() {
    cusparseHandle_t cusparseHandle;
    cusparseWarnIfError(cusparseCreate(&cusparseHandle), "cusparseCreate(&cusparseHandle)", __FILE__, __func__,
__LINE__);
}

Note

It is probably easier to use the macro OPM_CUSPARSE_WARN_IF_ERROR

Prefer the cusparseSafeCall/OPM_CUSPARSE_SAFE_CALL counterpart unless you really don't want to throw an exception.

Todo

Refactor to use std::source_location once we shift to C++20

extractNonzeroValues()

template<class T, class M>

std::vector< T > Opm::gpuistl::detail::extractNonzeroValues

(

const M &

matrix

)

Extract non-zero values from a DUNE matrix in the order expected by cuSPARSE.

This function iterates through the matrix and extracts the non-zero values in the Block Compressed Sparse Row (BSR) format order that cuSPARSE expects.

Template Parameters

T	the element type (float or double)
M	the matrix type (assumed to be DUNE::BCRSMatrix compatible)

Parameters

matrix

the input matrix

Returns

vector containing the non-zero elements in BSR order

extractSparsityPattern()

template<class M>

std::pair< std::vector< int >, std::vector< int > > Opm::gpuistl::detail::extractSparsityPattern

(

const M &

matrix

)

Extract sparsity pattern from a DUNE matrix.

This function extracts the row and column indices that define the sparsity pattern of a DUNE matrix in the CSR format expected by cuSPARSE.

Template Parameters

M	the matrix type (assumed to be DUNE::BCRSMatrix compatible)

Parameters

matrix

the input matrix

Returns

pair of (columnIndices, rowIndices) in CSR format

findFirstElementPointer()

template<class T, class M>

T * Opm::gpuistl::detail::findFirstElementPointer

(

const M &

matrix

)

Find pointer to first matrix element for direct memory access.

This function finds a pointer to the first element in the matrix data for cases where direct memory copying is desired.

Template Parameters

T	the element type
M	the matrix type

Parameters

matrix

the input matrix

Returns

pointer to first element, or nullptr if matrix is empty

getCublasErrorMessage()

std::string Opm::gpuistl::detail::getCublasErrorMessage ( cublasStatus_t error, const std::string_view & expression, const std::string_view & filename, const std::string_view & functionName, size_t lineNumber )

inline

getCublasErrorMessage generates the error message to display for a given error.

Parameters

error	the error code from cublas
expression	the expresison (say "cublasCreate(&handle)")
filename	the code file the error occured in (typically FILE)
functionName	name of the function the error occured in (typically func)
lineNumber	the line number the error occured in (typically LINE)

Todo

Refactor to use std::source_location once we shift to C++20

Returns

An error message to be displayed.

Note

This function is mostly for internal use.

getCudaErrorMessage()

std::string Opm::gpuistl::detail::getCudaErrorMessage ( cudaError_t error, const std::string_view & expression, const std::string_view & filename, const std::string_view & functionName, size_t lineNumber )

inline

getCudaErrorMessage generates the error message to display for a given error.

Parameters

error	the error code from cublas
expression	the expresison (say "cudaMalloc(&pointer, 1)")
filename	the code file the error occured in (typically FILE)
functionName	name of the function the error occured in (typically func)
lineNumber	the line number the error occured in (typically LINE)

Todo

Refactor to use std::source_location once we shift to C++20

Returns

An error message to be displayed.

Note

This function is mostly for internal use.

getCusparseErrorCodeToString()

std::string Opm::gpuistl::detail::getCusparseErrorCodeToString ( int code )

inline

getCusparseErrorCodeToString Converts an error code returned from a cusparse function a human readable string.

Parameters

code	an error code from a cusparse routine

Returns

a human readable string.

getCusparseErrorMessage()

std::string Opm::gpuistl::detail::getCusparseErrorMessage ( cusparseStatus_t error, const std::string_view & expression, const std::string_view & filename, const std::string_view & functionName, size_t lineNumber )

inline

getCusparseErrorMessage generates the error message to display for a given error.

Parameters

error	the error code from cublas
expression	the expresison (say "cusparseCreate(&handle)")
filename	the code file the error occured in (typically FILE)
functionName	name of the function the error occured in (typically func)
lineNumber	the line number the error occured in (typically LINE)

Todo

Refactor to use std::source_location once we shift to C++20

Returns

An error message to be displayed.

Note

This function is mostly for internal use.

getQuasiImpesWeights()

template<typename T, bool transpose>

void Opm::gpuistl::detail::getQuasiImpesWeights	(	const GpuSparseMatrixWrapper< T > &	matrix,
		std::size_t	pressureVarIndex,
		GpuVector< T > &	weights,
		const GpuVector< int > &	diagonalIndices )

Calculates quasi-IMPES weights for CPR preconditioner on GPU.

Parameters

matrix	The matrix used for weight calculation
pressureVarIndex	The index of the pressure variable
weights	Output vector to store the calculated weights
diagonalIndices	The diagonal indices of the matrix

Template Parameters

transpose

Whether to use the transpose mode for calculation

innerProductAtIndices()

template<class T>

T Opm::gpuistl::detail::innerProductAtIndices	(	cublasHandle_t	cublasHandle,
		const T *	deviceA,
		const T *	deviceB,
		T *	buffer,
		size_t	numberOfElements,
		const int *	indices )

innerProductAtIndices computes the inner product between deviceA[indices] and deviceB[indices]

Parameters

cublasHandle	a valid (initialized) cublas handle
deviceA	data A (device memory)
deviceB	data B (device memory)
buffer	a buffer with number of elements equal to numberOfElements (device memory)
numberOfElements	number of indices
indices	the indices to compute the inner product over (device memory)

Returns

the result of the inner product

Note

This is equivalent to projecting the vectors to the indices contained in indices, then doing the inner product of those projected vectors.

isGPUPointer() [1/3]

template<class T>

bool Opm::gpuistl::detail::isGPUPointer ( const std::shared_ptr< T > & ptr )

inline

Checks if the given std::shared_ptr refers to GPU memory.

Template Parameters

T	The type stored within the pointer.

Parameters

ptr	The std::shared_ptr object to inspect.

Returns

true if the pointer addresses GPU memory; otherwise false.

isGPUPointer() [2/3]

template<class T, class D>

bool Opm::gpuistl::detail::isGPUPointer ( const std::unique_ptr< T, D > & ptr )

inline

Checks if the given std::unique_ptr with custom deleter refers to GPU memory.

Template Parameters

T	The type stored within the pointer.
D	The custom deleter type.

Parameters

ptr	The std::unique_ptr object to inspect.

Returns

true if the pointer addresses GPU memory; otherwise false.

isGPUPointer() [3/3]

template<class T>

bool Opm::gpuistl::detail::isGPUPointer ( const T * ptr )

inline

Checks whether the given pointer is associated with GPU device memory.

This function retrieves CUDA pointer attributes for the provided pointer and determines whether it references device memory. It returns true if the pointer corresponds to GPU memory; otherwise, it returns false.

Template Parameters

T	Type of elements pointed to by the input pointer.

Parameters

ptr	Pointer to the memory that needs to be checked.

Returns

True if the pointer represents GPU memory, false otherwise.

makeCPUMatrix()

template<class Operator, class BlockType>

Dune::BCRSMatrix< BlockType > Opm::gpuistl::detail::makeCPUMatrix

(

const Operator &

op

)

Utility functions for GPU preconditioner creation.

These functions are temporary workarounds needed because some GPU preconditioners require CPU matrices for initial setup (e.g., graph coloring). They will be removed once GPU preconditioners have native GPU constructors. This function creates a CPU matrix from the operator holding a GPU matrix.

This is a workaround for now since some of the GPU preconditioners need a CPU matrix for the initial setup (graph coloring). The CPU matrix is only used in the constructor, not in the update function or the apply function.

makeMatrixWithNonzeroDiagonal()

template<class Matrix>

const Matrix Opm::gpuistl::detail::makeMatrixWithNonzeroDiagonal	(	const Matrix &	matrix,
		const typename Matrix::field_type	replacementValue = std::numeric_limits<typename Matrix::field_type>::epsilon() )

makeMatrixWithNonzeroDiagonal creates a new matrix with the zero diagonal elements (when viewed as a matrix of scalrars) set to replacementValue

Parameters

matrix	the matrix to replace
replacementValue	the value to set in the diagonal elements that are zero

Returns

a new matrix with non non-zero diagonal elements.

Note

This modification is needed for the CuSparse implementation of ILU0. While the the algorithm operates on block matrices, it still requires that the scalar matrix has no zero diagonal elements.

makeSafeMatrixDescriptor()

auto Opm::gpuistl::detail::makeSafeMatrixDescriptor ( )

inline

Create RAII-managed cuSPARSE sparse matrix descriptor.

Returns

unique_ptr with custom deleter for automatic cleanup

makeSafeVectorDescriptor()

auto Opm::gpuistl::detail::makeSafeVectorDescriptor ( )

inline

Create RAII-managed cuSPARSE dense vector descriptor.

Returns

unique_ptr with custom deleter for automatic cleanup

prolongateVector()

template<typename T, bool transpose>

void Opm::gpuistl::detail::prolongateVector	(	const GpuVector< T > &	coarse,
		GpuVector< T > &	fine,
		const GpuVector< T > &	weights,
		std::size_t	pressureVarIndex )

Prolongs a coarse level vector to a fine level vector based on pressure index.

Parameters

coarse	The coarse level vector
fine	The fine level vector to fill
weights	The weights vector for pressure calculation
pressureVarIndex	The index of the pressure variable

Template Parameters

transpose

Whether to use the transpose representation or not

restrictVector()

template<typename T, bool transpose>

void Opm::gpuistl::detail::restrictVector	(	const GpuVector< T > &	fine,
		GpuVector< T > &	coarse,
		const GpuVector< T > &	weights,
		std::size_t	pressureVarIndex )

Restricts a fine level vector to a coarse level vector based on pressure index.

Parameters

fine	The fine level vector
coarse	The coarse level vector to fill
weights	The weights vector for pressure calculation
pressureVarIndex	The index of the pressure variable

Template Parameters

transpose

Whether to use the transpose representation or not

setVectorValue()

template<class T>

void Opm::gpuistl::detail::setVectorValue	(	T *	deviceData,
		size_t	numberOfElements,
		const T &	value )

setVectorValue sets every element of deviceData to value

Parameters

deviceData	pointer to GPU memory
numberOfElements	number of elements to set to value
value	the value to use

setZeroAtIndexSet()

template<class T>

void Opm::gpuistl::detail::setZeroAtIndexSet	(	T *	deviceData,
		size_t	numberOfElements,
		const int *	indices )

setZeroAtIndexSet sets deviceData to zero in the indices of contained in indices

Parameters

deviceData	the data to operate on (device memory)
numberOfElements	number of indices
indices	the indices to use (device memory)

shouldCallPreconditionerPost()

template<class PreconditionerType>

bool Opm::gpuistl::detail::shouldCallPreconditionerPost ( )

constexpr

Tests (compile time) if the preconditioner type needs to call post() after a call to apply(...).

Note

This is mostly used to avoid unneeded copying back and front to the GPU, as well as avoiding communication.

shouldCallPreconditionerPre()

template<class PreconditionerType>

bool Opm::gpuistl::detail::shouldCallPreconditionerPre ( )

constexpr

Tests (compile time) if the preconditioner type needs to call pre() before a call to apply().

Note

This is mostly used to avoid unneeded copying back and front to the GPU, as well as avoiding communication.

to_int()

int Opm::gpuistl::detail::to_int ( std::size_t s )

inline

to_int converts a (on most relevant platforms) 64 bits unsigned size_t to a signed 32 bits signed int

Parameters

s	the unsigned integer

Exceptions

std::invalid_argument

exception if s is out of range for an int

Returns

converted s to int if s is within the range of int

Todo

This can be done for more generic types, but then it is probably wise to wait for C++20's cmp-functions

to_size_t()

__host__ __device__ std::size_t Opm::gpuistl::detail::to_size_t ( int i )

inline

to_size_t converts a (on most relevant platforms) a 32 bit signed int to a 64 bits unsigned int

Parameters

i	the signed integer

Returns

converted i to size_t if it is a non-negative integer.

Exceptions

std::invalid_argument

if i is negative.

Todo

This can be done for more generic types, but then it is probably wise to wait for C++20's cmp-functions

tuneThreadBlockSize()

template<typename func>

int Opm::gpuistl::detail::tuneThreadBlockSize	(	func &	f,
		std::string	descriptionOfFunction )

Function that tests the best thread block size, assumes the provided function depends on threadblock-size.

Template Parameters

The	type of the function to tune

Parameters

f	the function to tune, which takes the thread block size as the input
descriptionOfFunction	Description of function

validateMatrixCompatibility()

void Opm::gpuistl::detail::validateMatrixCompatibility ( size_t nonzeroes1, size_t blockSize1, size_t N1, size_t nonzeroes2, size_t blockSize2, size_t N2 )

inline

Validate that two matrices have compatible dimensions.

Parameters

nonzeroes1	number of nonzeros in first matrix
blockSize1	block size of first matrix
N1	number of rows in first matrix
nonzeroes2	number of nonzeros in second matrix
blockSize2	block size of second matrix
N2	number of rows in second matrix

validateMatrixConversion()

template<class M>

void Opm::gpuistl::detail::validateMatrixConversion	(	const M &	matrix,
		const std::vector< int > &	columnIndices,
		const std::vector< int > &	rowIndices )

Validate matrix conversion parameters.

Performs sanity checks on the extracted sparsity pattern to ensure it's valid for cuSPARSE operations.

Template Parameters

M	the matrix type

Parameters

matrix	the original matrix
columnIndices	the extracted column indices
rowIndices	the extracted row indices

validateVectorMatrixSizes()

void Opm::gpuistl::detail::validateVectorMatrixSizes ( size_t vectorSize, size_t matrixBlockSize, size_t matrixColumns )

inline

Validate vector-matrix size compatibility for operations.

Parameters

vectorSize	size of the vector
matrixBlockSize	block size of the matrix
matrixColumns	number of matrix columns

weightedDiagMV()

template<class T>

void Opm::gpuistl::detail::weightedDiagMV	(	const T *	squareBlockVector,
		const size_t	numberOfRows,
		const size_t	blocksize,
		T	relaxationFactor,
		const T *	srcVec,
		T *	dstVec )

Compue the weighted matrix vector product where the matrix is diagonal, the diagonal is a vector, meaning we compute the Hadamard product.

Parameters

	squareBlockVector	A GpuVector whose elements are NxN matrix blocks
	numberOfRows	The number of rows in the vector
	blocksize	The sidelength of the square block elements in the vector
	relaxationFactor	Relaxation factor to use
	srcVec	A pointer to the data of the GpuVector we multiply the blockvector with
[out]	dstVec	A pointer to the data of the GpuVector we store the result in

Note

This is implemented as a faster way to multiply a diagonal matrix with a blockvector. We need only store the diagonal of the matrix and use this product.