GpuSparseMatrix.hpp
Go to the documentation of this file.
1/*
2 Copyright 2022-2023 SINTEF AS
3
4 This file is part of the Open Porous Media project (OPM).
5
6 OPM is free software: you can redistribute it and/or modify
7 it under the terms of the GNU General Public License as published by
8 the Free Software Foundation, either version 3 of the License, or
9 (at your option) any later version.
10
11 OPM is distributed in the hope that it will be useful,
12 but WITHOUT ANY WARRANTY; without even the implied warranty of
13 MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
14 GNU General Public License for more details.
15
16 You should have received a copy of the GNU General Public License
17 along with OPM. If not, see <http://www.gnu.org/licenses/>.
18*/
19#ifndef OPM_GPUSPARSEMATRIX_HPP
20#define OPM_GPUSPARSEMATRIX_HPP
21#include <cusparse.h>
22#include <iostream>
23#include <memory>
24#include <opm/common/ErrorMacros.hpp>
29#include <vector>
30
31namespace Opm::gpuistl
32{
33
46template <typename T>
48{
49public:
63 GpuSparseMatrix(const T* nonZeroElements,
64 const int* rowIndices,
65 const int* columnIndices,
66 size_t numberOfNonzeroBlocks,
67 size_t blockSize,
68 size_t numberOfRows);
69
74
79
81
91 template <class MatrixType>
92 static GpuSparseMatrix<T> fromMatrix(const MatrixType& matrix, bool copyNonZeroElementsDirectly = false);
93
98
103
108
109
114
118 size_t N() const
119 {
120 // Technically this safe conversion is not needed since we enforce these to be
121 // non-negative in the constructor, but keeping them for added sanity for now.
122 //
123 // We don't believe this will yield any performance penality (it's used too far away from the inner loop),
124 // but should that be false, they can be removed.
125 return detail::to_size_t(m_numberOfRows);
126 }
127
132 size_t nonzeroes() const
133 {
134 // Technically this safe conversion is not needed since we enforce these to be
135 // non-negative in the constructor, but keeping them for added sanity for now.
136 //
137 // We don't believe this will yield any performance penality (it's used too far away from the inner loop),
138 // but should that be false, they can be removed.
139 return detail::to_size_t(m_numberOfNonzeroBlocks);
140 }
141
147 GpuVector<T>& getNonZeroValues()
148 {
149 return m_nonZeroElements;
150 }
151
157 const GpuVector<T>& getNonZeroValues() const
158 {
159 return m_nonZeroElements;
160 }
161
167 GpuVector<int>& getRowIndices()
168 {
169 return m_rowIndices;
170 }
171
177 const GpuVector<int>& getRowIndices() const
178 {
179 return m_rowIndices;
180 }
181
187 GpuVector<int>& getColumnIndices()
188 {
189 return m_columnIndices;
190 }
191
197 const GpuVector<int>& getColumnIndices() const
198 {
199 return m_columnIndices;
200 }
201
208 size_t dim() const
209 {
210 // Technically this safe conversion is not needed since we enforce these to be
211 // non-negative in the constructor, but keeping them for added sanity for now.
212 //
213 // We don't believe this will yield any performance penality (it's used too far away from the inner loop),
214 // but should that be false, they can be removed.
215 return detail::to_size_t(m_blockSize) * detail::to_size_t(m_numberOfRows);
216 }
217
221 size_t blockSize() const
222 {
223 // Technically this safe conversion is not needed since we enforce these to be
224 // non-negative in the constructor, but keeping them for added sanity for now.
225 //
226 // We don't believe this will yield any performance penality (it's used too far away from the inner loop),
227 // but should that be false, they can be removed.
228 return detail::to_size_t(m_blockSize);
229 }
230
237 {
238 return *m_matrixDescription;
239 }
240
248 virtual void mv(const GpuVector<T>& x, GpuVector<T>& y) const;
249
257 virtual void umv(const GpuVector<T>& x, GpuVector<T>& y) const;
258
259
267 virtual void usmv(T alpha, const GpuVector<T>& x, GpuVector<T>& y) const;
268
279 template <class MatrixType>
280 void updateNonzeroValues(const MatrixType& matrix, bool copyNonZeroElementsDirectly = false);
281
282private:
283 GpuVector<T> m_nonZeroElements;
284 GpuVector<int> m_columnIndices;
285 GpuVector<int> m_rowIndices;
286
287 // Notice that we store these three as int to make sure we are cusparse compatible.
288 //
289 // This gives the added benefit of checking the size constraints at construction of the matrix
290 // rather than in some call to cusparse.
291 const int m_numberOfNonzeroBlocks;
292 const int m_numberOfRows;
293 const int m_blockSize;
294
295 detail::GpuSparseMatrixDescriptionPtr m_matrixDescription;
296 detail::CuSparseHandle& m_cusparseHandle;
297
298 template <class VectorType>
299 void assertSameSize(const VectorType& vector) const;
300};
301} // namespace Opm::gpuistl
302#endif
The GpuSparseMatrix class simple wrapper class for a CuSparse matrix.
Definition: GpuSparseMatrix.hpp:48
GpuSparseMatrix & operator=(const GpuSparseMatrix &)=delete
virtual void mv(const GpuVector< T > &x, GpuVector< T > &y) const
mv performs matrix vector multiply y = Ax
const GpuVector< int > & getRowIndices() const
getRowIndices returns the row indices used to represent the BSR structure.
Definition: GpuSparseMatrix.hpp:177
const GpuVector< int > & getColumnIndices() const
getColumnIndices returns the column indices used to represent the BSR structure.
Definition: GpuSparseMatrix.hpp:197
void setNonUnitDiagonal()
setNonUnitDiagonal sets the CuSparse flag that this has non-unit diagional.
void updateNonzeroValues(const MatrixType &matrix, bool copyNonZeroElementsDirectly=false)
updateNonzeroValues updates the non-zero values by using the non-zero values of the supplied matrix
detail::GpuSparseMatrixDescription & getDescription()
getDescription the cusparse matrix description.
Definition: GpuSparseMatrix.hpp:236
GpuVector< int > & getRowIndices()
getRowIndices returns the row indices used to represent the BSR structure.
Definition: GpuSparseMatrix.hpp:167
const GpuVector< T > & getNonZeroValues() const
getNonZeroValues returns the GPU vector containing the non-zero values (ordered by block)
Definition: GpuSparseMatrix.hpp:157
size_t nonzeroes() const
nonzeroes behaves as the Dune::BCRSMatrix::nonzeros() function and returns the number of non zero blo...
Definition: GpuSparseMatrix.hpp:132
size_t blockSize() const
blockSize size of the blocks
Definition: GpuSparseMatrix.hpp:221
void setUpperTriangular()
setUpperTriangular sets the CuSparse flag that this is an upper diagonal (with unit diagonal) matrix.
size_t N() const
N returns the number of rows (which is equal to the number of columns)
Definition: GpuSparseMatrix.hpp:118
virtual void umv(const GpuVector< T > &x, GpuVector< T > &y) const
umv computes y=Ax+y
void setLowerTriangular()
setLowerTriangular sets the CuSparse flag that this is an lower diagonal (with non-unit diagonal) mat...
static GpuSparseMatrix< T > fromMatrix(const MatrixType &matrix, bool copyNonZeroElementsDirectly=false)
fromMatrix creates a new matrix with the same block size and values as the given matrix
void setUnitDiagonal()
setUnitDiagonal sets the CuSparse flag that this has unit diagional.
virtual void usmv(T alpha, const GpuVector< T > &x, GpuVector< T > &y) const
umv computes y=alpha * Ax + y
GpuVector< int > & getColumnIndices()
getColumnIndices returns the column indices used to represent the BSR structure.
Definition: GpuSparseMatrix.hpp:187
GpuSparseMatrix(const T *nonZeroElements, const int *rowIndices, const int *columnIndices, size_t numberOfNonzeroBlocks, size_t blockSize, size_t numberOfRows)
GpuSparseMatrix(const GpuSparseMatrix &)=delete
GpuVector< T > & getNonZeroValues()
getNonZeroValues returns the GPU vector containing the non-zero values (ordered by block)
Definition: GpuSparseMatrix.hpp:147
size_t dim() const
dim returns the dimension of the vector space on which this matrix acts
Definition: GpuSparseMatrix.hpp:208
The CuSparseHandle class provides a singleton for the simulator universal cuSparseHandle.
Definition: CuSparseHandle.hpp:41
The CuSparseResource class wraps a CuSparse resource in a proper RAII pattern.
Definition: CuSparseResource.hpp:55
std::shared_ptr< CuSparseResource< cusparseMatDescr_t > > GpuSparseMatrixDescriptionPtr
Definition: CuMatrixDescription.hpp:35
__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
Definition: safe_conversion.hpp:86
Definition: autotuner.hpp:29