1 /* Portions of this code are under:
2 Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
3 */
4 #pragma once
5
6 #include <petscpkg_version.h>
7 #include <../src/vec/vec/impls/seq/cupm/vecseqcupm.hpp> /* for VecSeq_CUPM */
8 #include <../src/sys/objects/device/impls/cupm/cupmthrustutility.hpp>
9 #include <petsc/private/petsclegacycupmblas.h>
10
11 #if PETSC_PKG_HIP_VERSION_GE(5, 2, 0)
12 #include <hipsparse/hipsparse.h>
13 #else /* PETSC_PKG_HIP_VERSION_GE(5,2,0) */
14 #include <hipsparse.h>
15 #endif /* PETSC_PKG_HIP_VERSION_GE(5,2,0) */
16 #include "hip/hip_runtime.h"
17
18 #include <algorithm>
19 #include <vector>
20
21 #include <thrust/device_vector.h>
22 #include <thrust/device_ptr.h>
23 #include <thrust/device_malloc_allocator.h>
24 #include <thrust/transform.h>
25 #include <thrust/functional.h>
26 #include <thrust/sequence.h>
27 #include <thrust/system/system_error.h>
28
29 #if defined(PETSC_USE_COMPLEX)
30 #if defined(PETSC_USE_REAL_SINGLE)
31 const hipComplex PETSC_HIPSPARSE_ONE = {1.0f, 0.0f};
32 const hipComplex PETSC_HIPSPARSE_ZERO = {0.0f, 0.0f};
33 #define hipsparseXcsrilu02_bufferSize(a, b, c, d, e, f, g, h, i) hipsparseCcsrilu02_bufferSize(a, b, c, d, (hipComplex *)e, f, g, h, i)
34 #define hipsparseXcsrilu02_analysis(a, b, c, d, e, f, g, h, i, j) hipsparseCcsrilu02_analysis(a, b, c, d, (hipComplex *)e, f, g, h, i, j)
35 #define hipsparseXcsrilu02(a, b, c, d, e, f, g, h, i, j) hipsparseCcsrilu02(a, b, c, d, (hipComplex *)e, f, g, h, i, j)
36 #define hipsparseXcsric02_bufferSize(a, b, c, d, e, f, g, h, i) hipsparseCcsric02_bufferSize(a, b, c, d, (hipComplex *)e, f, g, h, i)
37 #define hipsparseXcsric02_analysis(a, b, c, d, e, f, g, h, i, j) hipsparseCcsric02_analysis(a, b, c, d, (hipComplex *)e, f, g, h, i, j)
38 #define hipsparseXcsric02(a, b, c, d, e, f, g, h, i, j) hipsparseCcsric02(a, b, c, d, (hipComplex *)e, f, g, h, i, j)
39 #elif defined(PETSC_USE_REAL_DOUBLE)
40 const hipDoubleComplex PETSC_HIPSPARSE_ONE = {1.0, 0.0};
41 const hipDoubleComplex PETSC_HIPSPARSE_ZERO = {0.0, 0.0};
42 #define hipsparseXcsrilu02_bufferSize(a, b, c, d, e, f, g, h, i) hipsparseZcsrilu02_bufferSize(a, b, c, d, (hipDoubleComplex *)e, f, g, h, i)
43 #define hipsparseXcsrilu02_analysis(a, b, c, d, e, f, g, h, i, j) hipsparseZcsrilu02_analysis(a, b, c, d, (hipDoubleComplex *)e, f, g, h, i, j)
44 #define hipsparseXcsrilu02(a, b, c, d, e, f, g, h, i, j) hipsparseZcsrilu02(a, b, c, d, (hipDoubleComplex *)e, f, g, h, i, j)
45 #define hipsparseXcsric02_bufferSize(a, b, c, d, e, f, g, h, i) hipsparseZcsric02_bufferSize(a, b, c, d, (hipDoubleComplex *)e, f, g, h, i)
46 #define hipsparseXcsric02_analysis(a, b, c, d, e, f, g, h, i, j) hipsparseZcsric02_analysis(a, b, c, d, (hipDoubleComplex *)e, f, g, h, i, j)
47 #define hipsparseXcsric02(a, b, c, d, e, f, g, h, i, j) hipsparseZcsric02(a, b, c, d, (hipDoubleComplex *)e, f, g, h, i, j)
48 #endif /* Single or double */
49 #else /* not complex */
50 const PetscScalar PETSC_HIPSPARSE_ONE = 1.0;
51 const PetscScalar PETSC_HIPSPARSE_ZERO = 0.0;
52 #if defined(PETSC_USE_REAL_SINGLE)
53 #define hipsparseXcsrilu02_bufferSize hipsparseScsrilu02_bufferSize
54 #define hipsparseXcsrilu02_analysis hipsparseScsrilu02_analysis
55 #define hipsparseXcsrilu02 hipsparseScsrilu02
56 #define hipsparseXcsric02_bufferSize hipsparseScsric02_bufferSize
57 #define hipsparseXcsric02_analysis hipsparseScsric02_analysis
58 #define hipsparseXcsric02 hipsparseScsric02
59 #elif defined(PETSC_USE_REAL_DOUBLE)
60 #define hipsparseXcsrilu02_bufferSize hipsparseDcsrilu02_bufferSize
61 #define hipsparseXcsrilu02_analysis hipsparseDcsrilu02_analysis
62 #define hipsparseXcsrilu02 hipsparseDcsrilu02
63 #define hipsparseXcsric02_bufferSize hipsparseDcsric02_bufferSize
64 #define hipsparseXcsric02_analysis hipsparseDcsric02_analysis
65 #define hipsparseXcsric02 hipsparseDcsric02
66 #endif /* Single or double */
67 #endif /* complex or not */
68
69 #define csrsvInfo_t csrsv2Info_t
70 #define hipsparseCreateCsrsvInfo hipsparseCreateCsrsv2Info
71 #define hipsparseDestroyCsrsvInfo hipsparseDestroyCsrsv2Info
72 #if defined(PETSC_USE_COMPLEX)
73 #if defined(PETSC_USE_REAL_SINGLE)
74 #define hipsparseXcsrsv_buffsize(a, b, c, d, e, f, g, h, i, j) hipsparseCcsrsv2_bufferSize(a, b, c, d, e, (hipComplex *)(f), g, h, i, j)
75 #define hipsparseXcsrsv_analysis(a, b, c, d, e, f, g, h, i, j, k) hipsparseCcsrsv2_analysis(a, b, c, d, e, (const hipComplex *)(f), g, h, i, j, k)
76 #define hipsparseXcsrsv_solve(a, b, c, d, e, f, g, h, i, j, k, l, m, n) hipsparseCcsrsv2_solve(a, b, c, d, (const hipComplex *)(e), f, (const hipComplex *)(g), h, i, j, (const hipComplex *)(k), (hipComplex *)(l), m, n)
77 #elif defined(PETSC_USE_REAL_DOUBLE)
78 #define hipsparseXcsrsv_buffsize(a, b, c, d, e, f, g, h, i, j) hipsparseZcsrsv2_bufferSize(a, b, c, d, e, (hipDoubleComplex *)(f), g, h, i, j)
79 #define hipsparseXcsrsv_analysis(a, b, c, d, e, f, g, h, i, j, k) hipsparseZcsrsv2_analysis(a, b, c, d, e, (const hipDoubleComplex *)(f), g, h, i, j, k)
80 #define hipsparseXcsrsv_solve(a, b, c, d, e, f, g, h, i, j, k, l, m, n) hipsparseZcsrsv2_solve(a, b, c, d, (const hipDoubleComplex *)(e), f, (const hipDoubleComplex *)(g), h, i, j, (const hipDoubleComplex *)(k), (hipDoubleComplex *)(l), m, n)
81 #endif /* Single or double */
82 #else /* not complex */
83 #if defined(PETSC_USE_REAL_SINGLE)
84 #define hipsparseXcsrsv_buffsize hipsparseScsrsv2_bufferSize
85 #define hipsparseXcsrsv_analysis hipsparseScsrsv2_analysis
86 #define hipsparseXcsrsv_solve hipsparseScsrsv2_solve
87 #elif defined(PETSC_USE_REAL_DOUBLE)
88 #define hipsparseXcsrsv_buffsize hipsparseDcsrsv2_bufferSize
89 #define hipsparseXcsrsv_analysis hipsparseDcsrsv2_analysis
90 #define hipsparseXcsrsv_solve hipsparseDcsrsv2_solve
91 #endif /* Single or double */
92 #endif /* not complex */
93
94 #if PETSC_PKG_HIP_VERSION_GE(4, 5, 0)
95 // #define cusparse_csr2csc cusparseCsr2cscEx2
96 #if defined(PETSC_USE_COMPLEX)
97 #if defined(PETSC_USE_REAL_SINGLE)
98 #define hipsparse_scalartype HIP_C_32F
99 #define hipsparse_csr_spgeam(a, b, c, d, e, f, g, h, i, j, k, l, m, n, o, p, q, r, s, t) hipsparseCcsrgeam2(a, b, c, (hipComplex *)d, e, f, (hipComplex *)g, h, i, (hipComplex *)j, k, l, (hipComplex *)m, n, o, p, (hipComplex *)q, r, s, t)
100 #define hipsparse_csr_spgeam_bufferSize(a, b, c, d, e, f, g, h, i, j, k, l, m, n, o, p, q, r, s, t) \
101 hipsparseCcsrgeam2_bufferSizeExt(a, b, c, (hipComplex *)d, e, f, (hipComplex *)g, h, i, (hipComplex *)j, k, l, (hipComplex *)m, n, o, p, (hipComplex *)q, r, s, t)
102 #elif defined(PETSC_USE_REAL_DOUBLE)
103 #define hipsparse_scalartype HIP_C_64F
104 #define hipsparse_csr_spgeam(a, b, c, d, e, f, g, h, i, j, k, l, m, n, o, p, q, r, s, t) \
105 hipsparseZcsrgeam2(a, b, c, (hipDoubleComplex *)d, e, f, (hipDoubleComplex *)g, h, i, (hipDoubleComplex *)j, k, l, (hipDoubleComplex *)m, n, o, p, (hipDoubleComplex *)q, r, s, t)
106 #define hipsparse_csr_spgeam_bufferSize(a, b, c, d, e, f, g, h, i, j, k, l, m, n, o, p, q, r, s, t) \
107 hipsparseZcsrgeam2_bufferSizeExt(a, b, c, (hipDoubleComplex *)d, e, f, (hipDoubleComplex *)g, h, i, (hipDoubleComplex *)j, k, l, (hipDoubleComplex *)m, n, o, p, (hipDoubleComplex *)q, r, s, t)
108 #endif /* Single or double */
109 #else /* not complex */
110 #if defined(PETSC_USE_REAL_SINGLE)
111 #define hipsparse_scalartype HIP_R_32F
112 #define hipsparse_csr_spgeam hipsparseScsrgeam2
113 #define hipsparse_csr_spgeam_bufferSize hipsparseScsrgeam2_bufferSizeExt
114 #elif defined(PETSC_USE_REAL_DOUBLE)
115 #define hipsparse_scalartype HIP_R_64F
116 #define hipsparse_csr_spgeam hipsparseDcsrgeam2
117 #define hipsparse_csr_spgeam_bufferSize hipsparseDcsrgeam2_bufferSizeExt
118 #endif /* Single or double */
119 #endif /* complex or not */
120 #endif /* PETSC_PKG_HIP_VERSION_GE(4, 5, 0) */
121
122 #if defined(PETSC_USE_COMPLEX)
123 #if defined(PETSC_USE_REAL_SINGLE)
124 #define hipsparse_scalartype HIP_C_32F
125 #define hipsparse_csr_spmv(a, b, c, d, e, f, g, h, i, j, k, l, m) hipsparseCcsrmv((a), (b), (c), (d), (e), (hipComplex *)(f), (g), (hipComplex *)(h), (i), (j), (hipComplex *)(k), (hipComplex *)(l), (hipComplex *)(m))
126 #define hipsparse_csr_spmm(a, b, c, d, e, f, g, h, i, j, k, l, m, n, o, p) hipsparseCcsrmm((a), (b), (c), (d), (e), (f), (hipComplex *)(g), (h), (hipComplex *)(i), (j), (k), (hipComplex *)(l), (m), (hipComplex *)(n), (hipComplex *)(o), (p))
127 #define hipsparse_csr2csc(a, b, c, d, e, f, g, h, i, j, k, l) hipsparseCcsr2csc((a), (b), (c), (d), (hipComplex *)(e), (f), (g), (hipComplex *)(h), (i), (j), (k), (l))
128 #define hipsparse_hyb_spmv(a, b, c, d, e, f, g, h) hipsparseChybmv((a), (b), (hipComplex *)(c), (d), (e), (hipComplex *)(f), (hipComplex *)(g), (hipComplex *)(h))
129 #define hipsparse_csr2hyb(a, b, c, d, e, f, g, h, i, j) hipsparseCcsr2hyb((a), (b), (c), (d), (hipComplex *)(e), (f), (g), (h), (i), (j))
130 #define hipsparse_hyb2csr(a, b, c, d, e, f) hipsparseChyb2csr((a), (b), (c), (hipComplex *)(d), (e), (f))
131 #define hipsparse_csr_spgemm(a, b, c, d, e, f, g, h, i, j, k, l, m, n, o, p, q, r, s, t) hipsparseCcsrgemm(a, b, c, d, e, f, g, h, (hipComplex *)i, j, k, l, m, (hipComplex *)n, o, p, q, (hipComplex *)r, s, t)
132 // #define hipsparse_csr_spgeam(a, b, c, d, e, f, g, h, i, j, k, l, m, n, o, p, q, r, s) hipsparseCcsrgeam(a, b, c, (hipComplex *)d, e, f, (hipComplex *)g, h, i, (hipComplex *)j, k, l, (hipComplex *)m, n, o, p, (hipComplex *)q, r, s)
133 #elif defined(PETSC_USE_REAL_DOUBLE)
134 #define hipsparse_scalartype HIP_C_64F
135 #define hipsparse_csr_spmv(a, b, c, d, e, f, g, h, i, j, k, l, m) hipsparseZcsrmv((a), (b), (c), (d), (e), (hipDoubleComplex *)(f), (g), (hipDoubleComplex *)(h), (i), (j), (hipDoubleComplex *)(k), (hipDoubleComplex *)(l), (hipDoubleComplex *)(m))
136 #define hipsparse_csr_spmm(a, b, c, d, e, f, g, h, i, j, k, l, m, n, o, p) \
137 hipsparseZcsrmm((a), (b), (c), (d), (e), (f), (hipDoubleComplex *)(g), (h), (hipDoubleComplex *)(i), (j), (k), (hipDoubleComplex *)(l), (m), (hipDoubleComplex *)(n), (hipDoubleComplex *)(o), (p))
138 #define hipsparse_csr2csc(a, b, c, d, e, f, g, h, i, j, k, l) hipsparseZcsr2csc((a), (b), (c), (d), (hipDoubleComplex *)(e), (f), (g), (hipDoubleComplex *)(h), (i), (j), (k), (l))
139 #define hipsparse_hyb_spmv(a, b, c, d, e, f, g, h) hipsparseZhybmv((a), (b), (hipDoubleComplex *)(c), (d), (e), (hipDoubleComplex *)(f), (hipDoubleComplex *)(g), (hipDoubleComplex *)(h))
140 #define hipsparse_csr2hyb(a, b, c, d, e, f, g, h, i, j) hipsparseZcsr2hyb((a), (b), (c), (d), (hipDoubleComplex *)(e), (f), (g), (h), (i), (j))
141 #define hipsparse_hyb2csr(a, b, c, d, e, f) hipsparseZhyb2csr((a), (b), (c), (hipDoubleComplex *)(d), (e), (f))
142 #define hipsparse_csr_spgemm(a, b, c, d, e, f, g, h, i, j, k, l, m, n, o, p, q, r, s, t) hipsparseZcsrgemm(a, b, c, d, e, f, g, h, (hipDoubleComplex *)i, j, k, l, m, (hipDoubleComplex *)n, o, p, q, (hipDoubleComplex *)r, s, t)
143 // #define hipsparse_csr_spgeam(a, b, c, d, e, f, g, h, i, j, k, l, m, n, o, p, q, r, s) hipsparseZcsrgeam(a, b, c, (hipDoubleComplex *)d, e, f, (hipDoubleComplex *)g, h, i, (hipDoubleComplex *)j, k, l, (hipDoubleComplex *)m, n, o, p, (hipDoubleComplex *)q, r, s)
144 #endif /* Single or double */
145 #else /* not complex */
146 #if defined(PETSC_USE_REAL_SINGLE)
147 #define hipsparse_scalartype HIP_R_32F
148 #define hipsparse_csr_spmv hipsparseScsrmv
149 #define hipsparse_csr_spmm hipsparseScsrmm
150 #define hipsparse_csr2csc hipsparseScsr2csc
151 #define hipsparse_hyb_spmv hipsparseShybmv
152 #define hipsparse_csr2hyb hipsparseScsr2hyb
153 #define hipsparse_hyb2csr hipsparseShyb2csr
154 #define hipsparse_csr_spgemm hipsparseScsrgemm
155 // #define hipsparse_csr_spgeam hipsparseScsrgeam
156 #elif defined(PETSC_USE_REAL_DOUBLE)
157 #define hipsparse_scalartype HIP_R_64F
158 #define hipsparse_csr_spmv hipsparseDcsrmv
159 #define hipsparse_csr_spmm hipsparseDcsrmm
160 #define hipsparse_csr2csc hipsparseDcsr2csc
161 #define hipsparse_hyb_spmv hipsparseDhybmv
162 #define hipsparse_csr2hyb hipsparseDcsr2hyb
163 #define hipsparse_hyb2csr hipsparseDhyb2csr
164 #define hipsparse_csr_spgemm hipsparseDcsrgemm
165 // #define hipsparse_csr_spgeam hipsparseDcsrgeam
166 #endif /* Single or double */
167 #endif /* complex or not */
168
169 #define THRUSTINTARRAY32 thrust::device_vector<int>
170 #define THRUSTINTARRAY thrust::device_vector<PetscInt>
171 #define THRUSTARRAY thrust::device_vector<PetscScalar>
172
173 /* A CSR matrix nonzero structure */
174 struct CsrMatrix {
175 PetscInt num_rows;
176 PetscInt num_cols;
177 PetscInt num_entries;
178 THRUSTINTARRAY32 *row_offsets;
179 THRUSTINTARRAY32 *column_indices;
180 THRUSTARRAY *values;
181 };
182
183 /* This is struct holding the relevant data needed to a MatSolve */
184 struct Mat_SeqAIJHIPSPARSETriFactorStruct {
185 /* Data needed for triangular solve */
186 hipsparseMatDescr_t descr;
187 hipsparseOperation_t solveOp;
188 CsrMatrix *csrMat;
189 csrsvInfo_t solveInfo;
190 hipsparseSolvePolicy_t solvePolicy; /* whether level information is generated and used */
191 int solveBufferSize;
192 void *solveBuffer;
193 size_t csr2cscBufferSize; /* to transpose the triangular factor (only used for CUDA >= 11.0) */
194 void *csr2cscBuffer;
195 PetscScalar *AA_h; /* managed host buffer for moving values to the GPU */
196 };
197
198 /* This is a larger struct holding all the triangular factors for a solve, transpose solve, and any indices used in a reordering */
199 struct Mat_SeqAIJHIPSPARSETriFactors {
200 Mat_SeqAIJHIPSPARSETriFactorStruct *loTriFactorPtr; /* pointer for lower triangular (factored matrix) on GPU */
201 Mat_SeqAIJHIPSPARSETriFactorStruct *upTriFactorPtr; /* pointer for upper triangular (factored matrix) on GPU */
202 Mat_SeqAIJHIPSPARSETriFactorStruct *loTriFactorPtrTranspose; /* pointer for lower triangular (factored matrix) on GPU for the transpose (useful for BiCG) */
203 Mat_SeqAIJHIPSPARSETriFactorStruct *upTriFactorPtrTranspose; /* pointer for upper triangular (factored matrix) on GPU for the transpose (useful for BiCG)*/
204 THRUSTINTARRAY *rpermIndices; /* indices used for any reordering */
205 THRUSTINTARRAY *cpermIndices; /* indices used for any reordering */
206 THRUSTARRAY *workVector;
207 hipsparseHandle_t handle; /* a handle to the hipsparse library */
208 PetscInt nnz; /* number of nonzeros ... need this for accurate logging between ICC and ILU */
209 hipDeviceProp_t dev_prop;
210 PetscBool init_dev_prop;
211
212 /* csrilu0/csric0 appeared in earlier versions of AMD ROCm^{TM}, but we use it along with hipsparseSpSV,
213 which first appeared in hipsparse with ROCm-4.5.0.
214 */
215 #if PETSC_PKG_HIP_VERSION_GE(4, 5, 0)
216 PetscScalar *csrVal;
217 int *csrRowPtr, *csrColIdx; /* a,i,j of M. Using int since some hipsparse APIs only support 32-bit indices */
218
219 /* Mixed mat descriptor types? yes, different hipsparse APIs use different types */
220 hipsparseMatDescr_t matDescr_M;
221 hipsparseSpMatDescr_t spMatDescr_L, spMatDescr_U;
222 hipsparseSpSVDescr_t spsvDescr_L, spsvDescr_Lt, spsvDescr_U, spsvDescr_Ut;
223
224 hipsparseDnVecDescr_t dnVecDescr_X, dnVecDescr_Y;
225 PetscScalar *X, *Y; /* data array of dnVec X and Y */
226
227 /* Mixed size types? yes */
228 int factBufferSize_M; /* M ~= LU or LLt */
229 size_t spsvBufferSize_L, spsvBufferSize_Lt, spsvBufferSize_U, spsvBufferSize_Ut;
230 /* hipsparse needs various buffers for factorization and solve of L, U, Lt, or Ut.
231 To save memory, we share the factorization buffer with one of spsvBuffer_L/U.
232 */
233 void *factBuffer_M, *spsvBuffer_L, *spsvBuffer_U, *spsvBuffer_Lt, *spsvBuffer_Ut;
234
235 csrilu02Info_t ilu0Info_M;
236 csric02Info_t ic0Info_M;
237 int structural_zero, numerical_zero;
238 hipsparseSolvePolicy_t policy_M;
239
240 /* In MatSolveTranspose() for ILU0, we use the two flags to do on-demand solve */
241 PetscBool createdTransposeSpSVDescr; /* Have we created SpSV descriptors for Lt, Ut? */
242 PetscBool updatedTransposeSpSVAnalysis; /* Have we updated SpSV analysis with the latest L, U values? */
243
244 PetscLogDouble numericFactFlops; /* Estimated FLOPs in ILU0/ICC0 numeric factorization */
245 #endif
246 };
247
248 struct Mat_HipsparseSpMV {
249 PetscBool initialized; /* Don't rely on spmvBuffer != NULL to test if the struct is initialized, */
250 size_t spmvBufferSize; /* since I'm not sure if smvBuffer can be NULL even after hipsparseSpMV_bufferSize() */
251 void *spmvBuffer;
252 hipsparseDnVecDescr_t vecXDescr, vecYDescr; /* descriptor for the dense vectors in y=op(A)x */
253 };
254
255 /* This is struct holding the relevant data needed to a MatMult */
256 struct Mat_SeqAIJHIPSPARSEMultStruct {
257 void *mat; /* opaque pointer to a matrix. This could be either a hipsparseHybMat_t or a CsrMatrix */
258 hipsparseMatDescr_t descr; /* Data needed to describe the matrix for a multiply */
259 THRUSTINTARRAY *cprowIndices; /* compressed row indices used in the parallel SpMV */
260 PetscScalar *alpha_one; /* pointer to a device "scalar" storing the alpha parameter in the SpMV */
261 PetscScalar *beta_zero; /* pointer to a device "scalar" storing the beta parameter in the SpMV as zero*/
262 PetscScalar *beta_one; /* pointer to a device "scalar" storing the beta parameter in the SpMV as one */
263 hipsparseSpMatDescr_t matDescr; /* descriptor for the matrix, used by SpMV and SpMM */
264 Mat_HipsparseSpMV hipSpMV[3]; /* different Mat_CusparseSpMV structs for non-transpose, transpose, conj-transpose */
Mat_SeqAIJHIPSPARSEMultStructMat_SeqAIJHIPSPARSEMultStruct265 Mat_SeqAIJHIPSPARSEMultStruct() : matDescr(NULL)
266 {
267 for (int i = 0; i < 3; i++) hipSpMV[i].initialized = PETSC_FALSE;
268 }
269 };
270
271 /* This is a larger struct holding all the matrices for a SpMV, and SpMV Transpose */
272 struct Mat_SeqAIJHIPSPARSE {
273 Mat_SeqAIJHIPSPARSEMultStruct *mat; /* pointer to the matrix on the GPU */
274 Mat_SeqAIJHIPSPARSEMultStruct *matTranspose; /* pointer to the matrix on the GPU (for the transpose ... useful for BiCG) */
275 THRUSTARRAY *workVector; /* pointer to a workvector to which we can copy the relevant indices of a vector we want to multiply */
276 THRUSTINTARRAY32 *rowoffsets_gpu; /* rowoffsets on GPU in non-compressed-row format. It is used to convert CSR to CSC */
277 PetscInt nrows; /* number of rows of the matrix seen by GPU */
278 MatHIPSPARSEStorageFormat format; /* the storage format for the matrix on the device */
279 PetscBool use_cpu_solve; /* Use AIJ_Seq (I)LU solve */
280 hipStream_t stream; /* a stream for the parallel SpMV ... this is not owned and should not be deleted */
281 hipsparseHandle_t handle; /* a handle to the cusparse library ... this may not be owned (if we're working in parallel i.e. multiGPUs) */
282 PetscObjectState nonzerostate; /* track nonzero state to possibly recreate the GPU matrix */
283 size_t csr2cscBufferSize; /* stuff used to compute the matTranspose above */
284 void *csr2cscBuffer; /* This is used as a C struct and is calloc'ed by PetscNewLog() */
285 // hipsparseCsr2CscAlg_t csr2cscAlg; /* algorithms can be selected from command line options */
286 hipsparseSpMVAlg_t spmvAlg;
287 hipsparseSpMMAlg_t spmmAlg;
288 THRUSTINTARRAY *csr2csc_i;
289 THRUSTINTARRAY *coords; /* permutation array used in MatSeqAIJHIPSPARSEMergeMats */
290 };
291
292 typedef struct Mat_SeqAIJHIPSPARSETriFactors *Mat_SeqAIJHIPSPARSETriFactors_p;
293
294 PETSC_INTERN PetscErrorCode MatSeqAIJHIPSPARSECopyToGPU(Mat);
295 PETSC_INTERN PetscErrorCode MatSeqAIJHIPSPARSEMergeMats(Mat, Mat, MatReuse, Mat *);
296 PETSC_INTERN PetscErrorCode MatSeqAIJHIPSPARSETriFactors_Reset(Mat_SeqAIJHIPSPARSETriFactors_p *);
297
298 using VecSeq_HIP = Petsc::vec::cupm::impl::VecSeq_CUPM<Petsc::device::cupm::DeviceType::HIP>;
299
isHipMem(const void * data)300 static inline bool isHipMem(const void *data)
301 {
302 using namespace Petsc::device::cupm;
303 auto mtype = PETSC_MEMTYPE_HOST;
304
305 PetscFunctionBegin;
306 PetscCallAbort(PETSC_COMM_SELF, impl::Interface<DeviceType::HIP>::PetscCUPMGetMemType(data, &mtype));
307 PetscFunctionReturn(PetscMemTypeDevice(mtype));
308 }
309