110   Mat_SeqAIJHIPSPARSE *hipsparsestruct = (Mat_SeqAIJHIPSPARSE *)A->spptr;  in MatHIPSPARSESetFormat_SeqAIJHIPSPARSE()  local
115     hipsparsestruct->format = format;  in MatHIPSPARSESetFormat_SeqAIJHIPSPARSE()
118     hipsparsestruct->format = format;  in MatHIPSPARSESetFormat_SeqAIJHIPSPARSE()
152   Mat_SeqAIJHIPSPARSE *hipsparsestruct = (Mat_SeqAIJHIPSPARSE *)A->spptr;  in MatHIPSPARSESetUseCPUSolve_SeqAIJHIPSPARSE()  local
155   hipsparsestruct->use_cpu_solve = use_cpu;  in MatHIPSPARSESetUseCPUSolve_SeqAIJHIPSPARSE()
204   Mat_SeqAIJHIPSPARSE *hipsparsestruct = (Mat_SeqAIJHIPSPARSE *)B->spptr;  in MatLUFactorNumeric_SeqAIJHIPSPARSE()  local
213   if (!hipsparsestruct->use_cpu_solve) {  in MatLUFactorNumeric_SeqAIJHIPSPARSE()
226   if (!hipsparsestruct->use_cpu_solve) PetscCall(MatSeqAIJHIPSPARSEILUAnalysisAndCopyToGPU(B));  in MatLUFactorNumeric_SeqAIJHIPSPARSE()
234   Mat_SeqAIJHIPSPARSE      *hipsparsestruct = (Mat_SeqAIJHIPSPARSE *)A->spptr;  in MatSetFromOptions_SeqAIJHIPSPARSE()  local
239 …MV", "MatHIPSPARSESetFormat", MatHIPSPARSEStorageFormats, (PetscEnum)hipsparsestruct->format, (Pet…  in MatSetFromOptions_SeqAIJHIPSPARSE()
241 …ve", "MatHIPSPARSESetFormat", MatHIPSPARSEStorageFormats, (PetscEnum)hipsparsestruct->format, (Pet…  in MatSetFromOptions_SeqAIJHIPSPARSE()
243 …Use CPU (I)LU solve", "MatHIPSPARSESetUseCPUSolve", hipsparsestruct->use_cpu_solve, &hipsparsestru…  in MatSetFromOptions_SeqAIJHIPSPARSE()
244     if (flg) PetscCall(MatHIPSPARSESetUseCPUSolve(A, hipsparsestruct->use_cpu_solve));  in MatSetFromOptions_SeqAIJHIPSPARSE()
246 …MVAlg_t", MatHIPSPARSESpMVAlgorithms, (PetscEnum)hipsparsestruct->spmvAlg, (PetscEnum *)&hipsparse…  in MatSetFromOptions_SeqAIJHIPSPARSE()
250 …MMAlg_t", MatHIPSPARSESpMMAlgorithms, (PetscEnum)hipsparsestruct->spmmAlg, (PetscEnum *)&hipsparse…  in MatSetFromOptions_SeqAIJHIPSPARSE()
944   Mat_SeqAIJHIPSPARSE           *hipsparsestruct = (Mat_SeqAIJHIPSPARSE *)A->spptr;  in MatSeqAIJHIPSPARSEFormExplicitTranspose()  local
951   matstruct = (Mat_SeqAIJHIPSPARSEMultStruct *)hipsparsestruct->mat;  in MatSeqAIJHIPSPARSEFormExplicitTranspose()
953   matstructT = (Mat_SeqAIJHIPSPARSEMultStruct *)hipsparsestruct->matTranspose;  in MatSeqAIJHIPSPARSEFormExplicitTranspose()
958 …if (hipsparsestruct->format != MAT_HIPSPARSE_CSR) PetscCall(MatSeqAIJHIPSPARSEInvalidateTranspose(…  in MatSeqAIJHIPSPARSEFormExplicitTranspose()
959   if (!hipsparsestruct->matTranspose) { /* create hipsparse matrix */  in MatSeqAIJHIPSPARSEFormExplicitTranspose()
974     if (hipsparsestruct->format == MAT_HIPSPARSE_CSR) {  in MatSeqAIJHIPSPARSEFormExplicitTranspose()
984 …if (!hipsparsestruct->rowoffsets_gpu) hipsparsestruct->rowoffsets_gpu = new THRUSTINTARRAY32(A->rm…  in MatSeqAIJHIPSPARSEFormExplicitTranspose()
985       hipsparsestruct->rowoffsets_gpu->assign(a->i, a->i + A->rmap->n + 1);  in MatSeqAIJHIPSPARSEFormExplicitTranspose()
989 …} else if (hipsparsestruct->format == MAT_HIPSPARSE_ELL || hipsparsestruct->format == MAT_HIPSPARS…  in MatSeqAIJHIPSPARSEFormExplicitTranspose()
1000 …PetscCallHIPSPARSE(hipsparse_hyb2csr(hipsparsestruct->handle, matstruct->descr, (hipsparseHybMat_t…  in MatSeqAIJHIPSPARSEFormExplicitTranspose()
1010 …PetscCallHIPSPARSE(hipsparse_csr2csc(hipsparsestruct->handle, temp->num_rows, temp->num_cols, temp…  in MatSeqAIJHIPSPARSEFormExplicitTranspose()
1016 …hipsparseHybPartition_t partition = hipsparsestruct->format == MAT_HIPSPARSE_ELL ? HIPSPARSE_HYB_P…  in MatSeqAIJHIPSPARSEFormExplicitTranspose()
1017 …PetscCallHIPSPARSE(hipsparse_csr2hyb(hipsparsestruct->handle, A->rmap->n, A->cmap->n, matstructT->…  in MatSeqAIJHIPSPARSEFormExplicitTranspose()
1037 …if (hipsparsestruct->format == MAT_HIPSPARSE_CSR) { /* transpose mat struct may be already present…  in MatSeqAIJHIPSPARSEFormExplicitTranspose()
1048 …if (!hipsparsestruct->rowoffsets_gpu) { /* this may be absent when we did not construct the transp…  in MatSeqAIJHIPSPARSEFormExplicitTranspose()
1049       hipsparsestruct->rowoffsets_gpu = new THRUSTINTARRAY32(A->rmap->n + 1);  in MatSeqAIJHIPSPARSEFormExplicitTranspose()
1050       hipsparsestruct->rowoffsets_gpu->assign(a->i, a->i + A->rmap->n + 1);  in MatSeqAIJHIPSPARSEFormExplicitTranspose()
1053     if (!hipsparsestruct->csr2csc_i) {  in MatSeqAIJHIPSPARSEFormExplicitTranspose()
1062 …RSE(hipsparse_csr2csc(hipsparsestruct->handle, A->rmap->n, A->cmap->n, matrix->num_entries, csr2cs…  in MatSeqAIJHIPSPARSEFormExplicitTranspose()
1068       hipsparsestruct->csr2csc_i = new THRUSTINTARRAY(matrix->num_entries);  in MatSeqAIJHIPSPARSEFormExplicitTranspose()
1069 …rm(thrust::device, matrixT->values->begin(), matrixT->values->end(), hipsparsestruct->csr2csc_i->b…  in MatSeqAIJHIPSPARSEFormExplicitTranspose()
1072 …x->values->begin(), hipsparsestruct->csr2csc_i->begin()), thrust::make_permutation_iterator(matrix…  in MatSeqAIJHIPSPARSEFormExplicitTranspose()
2013   Mat_SeqAIJHIPSPARSE           *hipsparsestruct = (Mat_SeqAIJHIPSPARSE *)A->spptr;  in MatSeqAIJHIPSPARSECopyToGPU()  local
2014   Mat_SeqAIJHIPSPARSEMultStruct *matstruct       = hipsparsestruct->mat;  in MatSeqAIJHIPSPARSECopyToGPU()
2022 …if (A->nonzerostate == hipsparsestruct->nonzerostate && hipsparsestruct->format == MAT_HIPSPARSE_C…  in MatSeqAIJHIPSPARSECopyToGPU()
2024       matrix = (CsrMatrix *)hipsparsestruct->mat->mat;  in MatSeqAIJHIPSPARSECopyToGPU()
2036 …   PetscCall(MatSeqAIJHIPSPARSEMultStruct_Destroy(&hipsparsestruct->mat, hipsparsestruct->format));  in MatSeqAIJHIPSPARSECopyToGPU()
2038       delete hipsparsestruct->workVector;  in MatSeqAIJHIPSPARSECopyToGPU()
2039       delete hipsparsestruct->rowoffsets_gpu;  in MatSeqAIJHIPSPARSECopyToGPU()
2040       hipsparsestruct->workVector     = NULL;  in MatSeqAIJHIPSPARSECopyToGPU()
2041       hipsparsestruct->rowoffsets_gpu = NULL;  in MatSeqAIJHIPSPARSECopyToGPU()
2060         hipsparsestruct->nrows = m;  in MatSeqAIJHIPSPARSECopyToGPU()
2072 …PetscCallHIPSPARSE(hipsparseSetPointerMode(hipsparsestruct->handle, HIPSPARSE_POINTER_MODE_DEVICE)…  in MatSeqAIJHIPSPARSECopyToGPU()
2075         if (hipsparsestruct->format == MAT_HIPSPARSE_CSR) {  in MatSeqAIJHIPSPARSECopyToGPU()
2094 …} else if (hipsparsestruct->format == MAT_HIPSPARSE_ELL || hipsparsestruct->format == MAT_HIPSPARS…  in MatSeqAIJHIPSPARSECopyToGPU()
2108 …hipsparseHybPartition_t partition = hipsparsestruct->format == MAT_HIPSPARSE_ELL ? HIPSPARSE_HYB_P…  in MatSeqAIJHIPSPARSECopyToGPU()
2109 …PetscCallHIPSPARSE(hipsparse_csr2hyb(hipsparsestruct->handle, mat->num_rows, mat->num_cols, matstr…  in MatSeqAIJHIPSPARSECopyToGPU()
2123           hipsparsestruct->workVector = new THRUSTARRAY(m);  in MatSeqAIJHIPSPARSECopyToGPU()
2128           hipsparsestruct->workVector = NULL;  in MatSeqAIJHIPSPARSECopyToGPU()
2135         hipsparsestruct->mat = matstruct;  in MatSeqAIJHIPSPARSECopyToGPU()
2141       hipsparsestruct->nonzerostate = A->nonzerostate;  in MatSeqAIJHIPSPARSECopyToGPU()
3024   Mat_SeqAIJHIPSPARSE           *hipsparsestruct = (Mat_SeqAIJHIPSPARSE *)A->spptr;  in MatMultAddKernel_SeqAIJHIPSPARSE()  local
3041     matstruct = (Mat_SeqAIJHIPSPARSEMultStruct *)hipsparsestruct->mat;  in MatMultAddKernel_SeqAIJHIPSPARSE()
3046       matstruct = (Mat_SeqAIJHIPSPARSEMultStruct *)hipsparsestruct->mat;  in MatMultAddKernel_SeqAIJHIPSPARSE()
3048       if (!hipsparsestruct->matTranspose) PetscCall(MatSeqAIJHIPSPARSEFormExplicitTranspose(A));  in MatMultAddKernel_SeqAIJHIPSPARSE()
3049       matstruct = (Mat_SeqAIJHIPSPARSEMultStruct *)hipsparsestruct->matTranspose;  in MatMultAddKernel_SeqAIJHIPSPARSE()
3066       dptr = compressed ? hipsparsestruct->workVector->data().get() : zarray;  in MatMultAddKernel_SeqAIJHIPSPARSE()
3071       if (hipsparsestruct->format == MAT_HIPSPARSE_CSR) {  in MatMultAddKernel_SeqAIJHIPSPARSE()
3081       xptr = compressed ? hipsparsestruct->workVector->data().get() : xarray;  in MatMultAddKernel_SeqAIJHIPSPARSE()
3090 …thrust::make_zip_iterator(thrust::make_tuple(hipsparsestruct->workVector->begin(), thrust::make_pe…  in MatMultAddKernel_SeqAIJHIPSPARSE()
3091 …thrust::make_zip_iterator(thrust::make_tuple(hipsparsestruct->workVector->begin(), thrust::make_pe…  in MatMultAddKernel_SeqAIJHIPSPARSE()
3093       if (hipsparsestruct->format == MAT_HIPSPARSE_CSR) {  in MatMultAddKernel_SeqAIJHIPSPARSE()
3100     if (hipsparsestruct->format == MAT_HIPSPARSE_CSR) {  in MatMultAddKernel_SeqAIJHIPSPARSE()
3106 …hipsparsestruct->handle, opA, matstruct->alpha_one, matstruct->matDescr, matstruct->hipSpMV[opA].v…  in MatMultAddKernel_SeqAIJHIPSPARSE()
3115 …PetscCallHIPSPARSE(hipsparseSpMV(hipsparsestruct->handle, opA, matstruct->alpha_one, matstruct->ma…  in MatMultAddKernel_SeqAIJHIPSPARSE()
3116 …escr, beta, matstruct->hipSpMV[opA].vecYDescr, hipsparse_scalartype, hipsparsestruct->spmvAlg, mat…  in MatMultAddKernel_SeqAIJHIPSPARSE()
3121 …PetscCallHIPSPARSE(hipsparse_csr_spmv(hipsparsestruct->handle, opA, nx, ny, mat->num_entries, mats…  in MatMultAddKernel_SeqAIJHIPSPARSE()
3124       if (hipsparsestruct->nrows) {  in MatMultAddKernel_SeqAIJHIPSPARSE()
3126 …PetscCallHIPSPARSE(hipsparse_hyb_spmv(hipsparsestruct->handle, opA, matstruct->alpha_one, matstruc…  in MatMultAddKernel_SeqAIJHIPSPARSE()
3151         thrust::async::for_each(thrust::hip::par.on(hipsparsestruct->stream),  in MatMultAddKernel_SeqAIJHIPSPARSE()
3152 …thrust::make_zip_iterator(thrust::make_tuple(hipsparsestruct->workVector->begin(), thrust::make_pe…  in MatMultAddKernel_SeqAIJHIPSPARSE()
3153 …thrust::make_zip_iterator(thrust::make_tuple(hipsparsestruct->workVector->begin(), thrust::make_pe…  in MatMultAddKernel_SeqAIJHIPSPARSE()
3157 … 0, PetscDefaultHipStream, n, matstruct->cprowIndices->data().get(), hipsparsestruct->workVector->…  in MatMultAddKernel_SeqAIJHIPSPARSE()