opm-simulators
GpuSparseMatrix.hpp
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 
22 #include <opm/common/ErrorMacros.hpp>
23 #include <opm/simulators/linalg/gpuistl/detail/CuMatrixDescription.hpp>
24 #include <opm/simulators/linalg/gpuistl/detail/CuSparseHandle.hpp>
25 #include <opm/simulators/linalg/gpuistl/detail/safe_conversion.hpp>
26 #include <opm/simulators/linalg/gpuistl/GpuVector.hpp>
27 #include <opm/simulators/linalg/gpuistl/GpuSparseMatrixGeneric.hpp>
28 
29 #include <cstddef>
30 #include <cusparse.h>
31 #include <memory>
32 #include <stdexcept>
33 #include <type_traits>
34 
35 namespace Opm::gpuistl
36 {
37 
58 template <typename T>
60 
61 {
62 public:
63  using field_type = T;
64 
73  static constexpr int max_block_size = 6;
74 
88  GpuSparseMatrix(const T* nonZeroElements,
89  const int* rowIndices,
90  const int* columnIndices,
91  std::size_t numberOfNonzeroBlocks,
92  std::size_t blockSize,
93  std::size_t numberOfRows);
94 
105  GpuSparseMatrix(const GpuVector<int>& rowIndices,
106  const GpuVector<int>& columnIndices,
107  std::size_t blockSize);
108 
110 
111  // We want to have this as non-mutable as possible, that is we do not want
112  // to deal with changing matrix sizes and sparsity patterns.
113  GpuSparseMatrix& operator=(const GpuSparseMatrix&) = delete;
114 
115  virtual ~GpuSparseMatrix();
116 
126  template <class MatrixType>
127  static GpuSparseMatrix<T> fromMatrix(const MatrixType& matrix, bool copyNonZeroElementsDirectly = false);
128 
132  void setUpperTriangular();
133 
137  void setLowerTriangular();
138 
142  void setUnitDiagonal();
143 
144 
148  void setNonUnitDiagonal();
149 
153  std::size_t N() const
154  {
155  // Technically this safe conversion is not needed since we enforce these to be
156  // non-negative in the constructor, but keeping them for added sanity for now.
157  //
158  // We don't believe this will yield any performance penality (it's used too far away from the inner loop),
159  // but should that be false, they can be removed.
160  return detail::to_size_t(m_numberOfRows);
161  }
162 
167  std::size_t nonzeroes() const
168  {
169  // Technically this safe conversion is not needed since we enforce these to be
170  // non-negative in the constructor, but keeping them for added sanity for now.
171  //
172  // We don't believe this will yield any performance penality (it's used too far away from the inner loop),
173  // but should that be false, they can be removed.
174  return detail::to_size_t(m_numberOfNonzeroBlocks);
175  }
176 
183  {
184  if (m_genericMatrixForBlockSize1) {
185  return m_genericMatrixForBlockSize1->getNonZeroValues();
186  }
187  return m_nonZeroElements;
188  }
189 
196  {
197  if (m_genericMatrixForBlockSize1) {
198  return m_genericMatrixForBlockSize1->getNonZeroValues();
199  }
200  return m_nonZeroElements;
201  }
202 
209  {
210  if (m_genericMatrixForBlockSize1) {
211  return m_genericMatrixForBlockSize1->getRowIndices();
212  }
213  return m_rowIndices;
214  }
215 
222  {
223  if (m_genericMatrixForBlockSize1) {
224  return m_genericMatrixForBlockSize1->getRowIndices();
225  }
226  return m_rowIndices;
227  }
228 
235  {
236  if (m_genericMatrixForBlockSize1) {
237  return m_genericMatrixForBlockSize1->getColumnIndices();
238  }
239  return m_columnIndices;
240  }
241 
248  {
249  if (m_genericMatrixForBlockSize1) {
250  return m_genericMatrixForBlockSize1->getColumnIndices();
251  }
252  return m_columnIndices;
253  }
254 
261  std::size_t dim() const
262  {
263  // Technically this safe conversion is not needed since we enforce these to be
264  // non-negative in the constructor, but keeping them for added sanity for now.
265  //
266  // We don't believe this will yield any performance penality (it's used too far away from the inner loop),
267  // but should that be false, they can be removed.
268  return detail::to_size_t(m_blockSize) * detail::to_size_t(m_numberOfRows);
269  }
270 
274  std::size_t blockSize() const
275  {
276  // Technically this safe conversion is not needed since we enforce these to be
277  // non-negative in the constructor, but keeping them for added sanity for now.
278  //
279  // We don't believe this will yield any performance penality (it's used too far away from the inner loop),
280  // but should that be false, they can be removed.
281  return detail::to_size_t(m_blockSize);
282  }
283 
290  {
291  return *m_matrixDescription;
292  }
293 
299  virtual void mv(const GpuVector<T>& x, GpuVector<T>& y) const;
300 
306  virtual void umv(const GpuVector<T>& x, GpuVector<T>& y) const;
307 
308 
315  virtual void usmv(T alpha, const GpuVector<T>& x, GpuVector<T>& y) const;
316 
327  template <class MatrixType>
328  void updateNonzeroValues(const MatrixType& matrix, bool copyNonZeroElementsDirectly = false);
329 
335  void updateNonzeroValues(const GpuSparseMatrix<T>& matrix);
336 
343 
347  void setToZero();
348 
369  template<class FunctionType>
370  auto dispatchOnBlocksize(FunctionType function) const
371  {
372  return dispatchOnBlocksizeImpl<max_block_size>(function);
373  }
374 
375 private:
376  GpuVector<T> m_nonZeroElements;
377  GpuVector<int> m_columnIndices;
378  GpuVector<int> m_rowIndices;
379 
380  // Notice that we store these three as int to make sure we are cusparse compatible.
381  //
382  // This gives the added benefit of checking the size constraints at construction of the matrix
383  // rather than in some call to cusparse.
384  const int m_numberOfNonzeroBlocks;
385  const int m_numberOfRows;
386  const int m_blockSize;
387 
388  detail::GpuSparseMatrixDescriptionPtr m_matrixDescription;
389  detail::CuSparseHandle& m_cusparseHandle;
390 
391  // For blockSize == 1, we use the generic API
392  std::unique_ptr<GpuSparseMatrixGeneric<T>> m_genericMatrixForBlockSize1;
393 
394  template <class VectorType>
395  void assertSameSize(const VectorType& vector) const;
396 
397  template<int blockSizeCompileTime, class FunctionType>
398  auto dispatchOnBlocksizeImpl(FunctionType function) const
399  {
400  if (blockSizeCompileTime == m_blockSize) {
401  return function(std::integral_constant<int, blockSizeCompileTime>());
402  }
403 
404  if constexpr (blockSizeCompileTime > 1) {
405  return dispatchOnBlocksizeImpl<blockSizeCompileTime - 1>(function);
406  } else {
407  OPM_THROW(std::runtime_error, fmt::format(fmt::runtime("Unsupported block size: {}"), m_blockSize));
408  }
409  }
410 };
411 
412 } // namespace Opm::gpuistl
413 
414 #endif
The GpuSparseMatrixGeneric class uses cuSPARSE Generic API for sparse matrix operations.
Definition: gpu_type_detection.hpp:34
const GpuVector< T > & getNonZeroValues() const
getNonZeroValues returns the GPU vector containing the non-zero values (ordered by block) ...
Definition: GpuSparseMatrix.hpp:195
virtual void umv(const GpuVector< T > &x, GpuVector< T > &y) const
umv computes y=Ax+y
Definition: GpuSparseMatrix.cpp:283
const GpuVector< int > & getColumnIndices() const
getColumnIndices returns the column indices used to represent the BSR structure.
Definition: GpuSparseMatrix.hpp:247
auto dispatchOnBlocksize(FunctionType function) const
Dispatches a function based on the block size of the matrix.
Definition: GpuSparseMatrix.hpp:370
The CuSparseHandle class provides a singleton for the simulator universal cuSparseHandle.
Definition: CuSparseHandle.hpp:40
void setUnitDiagonal()
setUnitDiagonal sets the CuSparse flag that this has unit diagional.
Definition: GpuSparseMatrix.cpp:234
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 ...
Definition: GpuSparseMatrix.cpp:116
void setToZero()
setToZero resets the matrix to zero values.
Definition: GpuSparseMatrix.cpp:207
GpuVector< int > & getRowIndices()
getRowIndices returns the row indices used to represent the BSR structure.
Definition: GpuSparseMatrix.hpp:208
void updateNonzeroValues(const MatrixType &matrix, bool copyNonZeroElementsDirectly=false)
updateNonzeroValues updates the non-zero values by using the non-zero values of the supplied matrix ...
Definition: GpuSparseMatrix.cpp:149
void setLowerTriangular()
setLowerTriangular sets the CuSparse flag that this is an lower diagonal (with non-unit diagonal) mat...
Definition: GpuSparseMatrix.cpp:227
__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:97
std::size_t blockSize() const
blockSize size of the blocks
Definition: GpuSparseMatrix.hpp:274
virtual void usmv(T alpha, const GpuVector< T > &x, GpuVector< T > &y) const
umv computes y=alpha * Ax + y
Definition: GpuSparseMatrix.cpp:319
GpuVector< int > & getColumnIndices()
getColumnIndices returns the column indices used to represent the BSR structure.
Definition: GpuSparseMatrix.hpp:234
GpuVector< T > & getNonZeroValues()
getNonZeroValues returns the GPU vector containing the non-zero values (ordered by block) ...
Definition: GpuSparseMatrix.hpp:182
GpuSparseMatrix(const T *nonZeroElements, const int *rowIndices, const int *columnIndices, std::size_t numberOfNonzeroBlocks, std::size_t blockSize, std::size_t numberOfRows)
Create the sparse matrix specified by the raw data.
std::shared_ptr< CuSparseResource< cusparseMatDescr_t > > GpuSparseMatrixDescriptionPtr
Pointer to GpuSparseMatrixDescription holder.
Definition: CuMatrixDescription.hpp:35
A small, fixed‑dimension MiniVector class backed by std::array that can be used in both host and CUD...
Definition: AmgxInterface.hpp:37
std::size_t dim() const
dim returns the dimension of the vector space on which this matrix acts
Definition: GpuSparseMatrix.hpp:261
std::size_t nonzeroes() const
nonzeroes behaves as the Dune::BCRSMatrix::nonzeros() function and returns the number of non zero blo...
Definition: GpuSparseMatrix.hpp:167
std::size_t N() const
N returns the number of rows (which is equal to the number of columns)
Definition: GpuSparseMatrix.hpp:153
The CuSparseResource class wraps a CuSparse resource in a proper RAII pattern.
Definition: CuSparseResource.hpp:54
static constexpr int max_block_size
Maximum block size supported by this implementation.
Definition: GpuSparseMatrix.hpp:73
The GpuSparseMatrix class simple wrapper class for a CuSparse matrix.
Definition: GpuSparseMatrix.hpp:59
detail::GpuSparseMatrixDescription & getDescription()
getDescription the cusparse matrix description.
Definition: GpuSparseMatrix.hpp:289
void setNonUnitDiagonal()
setNonUnitDiagonal sets the CuSparse flag that this has non-unit diagional.
Definition: GpuSparseMatrix.cpp:241
void setUpperTriangular()
setUpperTriangular sets the CuSparse flag that this is an upper diagonal (with unit diagonal) matrix...
Definition: GpuSparseMatrix.cpp:220
const GpuVector< int > & getRowIndices() const
getRowIndices returns the row indices used to represent the BSR structure.
Definition: GpuSparseMatrix.hpp:221
virtual void mv(const GpuVector< T > &x, GpuVector< T > &y) const
mv performs matrix vector multiply y = Ax
Definition: GpuSparseMatrix.cpp:248