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