23   MatCOOStruct_MPIAIJ *coo = *(MatCOOStruct_MPIAIJ **)data;  in MatCOOStructDestroy_MPIAIJHIPSPARSE()  local
26   PetscCall(PetscSFDestroy(&coo->sf));  in MatCOOStructDestroy_MPIAIJHIPSPARSE()
27   PetscCallHIP(hipFree(coo->Ajmap1));  in MatCOOStructDestroy_MPIAIJHIPSPARSE()
28   PetscCallHIP(hipFree(coo->Aperm1));  in MatCOOStructDestroy_MPIAIJHIPSPARSE()
29   PetscCallHIP(hipFree(coo->Bjmap1));  in MatCOOStructDestroy_MPIAIJHIPSPARSE()
30   PetscCallHIP(hipFree(coo->Bperm1));  in MatCOOStructDestroy_MPIAIJHIPSPARSE()
31   PetscCallHIP(hipFree(coo->Aimap2));  in MatCOOStructDestroy_MPIAIJHIPSPARSE()
32   PetscCallHIP(hipFree(coo->Ajmap2));  in MatCOOStructDestroy_MPIAIJHIPSPARSE()
33   PetscCallHIP(hipFree(coo->Aperm2));  in MatCOOStructDestroy_MPIAIJHIPSPARSE()
34   PetscCallHIP(hipFree(coo->Bimap2));  in MatCOOStructDestroy_MPIAIJHIPSPARSE()
35   PetscCallHIP(hipFree(coo->Bjmap2));  in MatCOOStructDestroy_MPIAIJHIPSPARSE()
36   PetscCallHIP(hipFree(coo->Bperm2));  in MatCOOStructDestroy_MPIAIJHIPSPARSE()
37   PetscCallHIP(hipFree(coo->Cperm1));  in MatCOOStructDestroy_MPIAIJHIPSPARSE()
38   PetscCallHIP(hipFree(coo->sendbuf));  in MatCOOStructDestroy_MPIAIJHIPSPARSE()
39   PetscCallHIP(hipFree(coo->recvbuf));  in MatCOOStructDestroy_MPIAIJHIPSPARSE()
40   PetscCall(PetscFree(coo));  in MatCOOStructDestroy_MPIAIJHIPSPARSE()
166   MatCOOStruct_MPIAIJ *coo;  in MatSetValuesCOO_MPIAIJHIPSPARSE()  local
171   PetscCall(PetscContainerGetPointer(container, &coo));  in MatSetValuesCOO_MPIAIJHIPSPARSE()
173   const auto &Annz   = coo->Annz;  in MatSetValuesCOO_MPIAIJHIPSPARSE()
174   const auto &Annz2  = coo->Annz2;  in MatSetValuesCOO_MPIAIJHIPSPARSE()
175   const auto &Bnnz   = coo->Bnnz;  in MatSetValuesCOO_MPIAIJHIPSPARSE()
176   const auto &Bnnz2  = coo->Bnnz2;  in MatSetValuesCOO_MPIAIJHIPSPARSE()
177   const auto &vsend  = coo->sendbuf;  in MatSetValuesCOO_MPIAIJHIPSPARSE()
178   const auto &v2     = coo->recvbuf;  in MatSetValuesCOO_MPIAIJHIPSPARSE()
179   const auto &Ajmap1 = coo->Ajmap1;  in MatSetValuesCOO_MPIAIJHIPSPARSE()
180   const auto &Ajmap2 = coo->Ajmap2;  in MatSetValuesCOO_MPIAIJHIPSPARSE()
181   const auto &Aimap2 = coo->Aimap2;  in MatSetValuesCOO_MPIAIJHIPSPARSE()
182   const auto &Bjmap1 = coo->Bjmap1;  in MatSetValuesCOO_MPIAIJHIPSPARSE()
183   const auto &Bjmap2 = coo->Bjmap2;  in MatSetValuesCOO_MPIAIJHIPSPARSE()
184   const auto &Bimap2 = coo->Bimap2;  in MatSetValuesCOO_MPIAIJHIPSPARSE()
185   const auto &Aperm1 = coo->Aperm1;  in MatSetValuesCOO_MPIAIJHIPSPARSE()
186   const auto &Aperm2 = coo->Aperm2;  in MatSetValuesCOO_MPIAIJHIPSPARSE()
187   const auto &Bperm1 = coo->Bperm1;  in MatSetValuesCOO_MPIAIJHIPSPARSE()
188   const auto &Bperm2 = coo->Bperm2;  in MatSetValuesCOO_MPIAIJHIPSPARSE()
189   const auto &Cperm1 = coo->Cperm1;  in MatSetValuesCOO_MPIAIJHIPSPARSE()
193     PetscCallHIP(hipMalloc((void **)&v1, coo->n * sizeof(PetscScalar)));  in MatSetValuesCOO_MPIAIJHIPSPARSE()
194     PetscCallHIP(hipMemcpy((void *)v1, v, coo->n * sizeof(PetscScalar), hipMemcpyHostToDevice));  in MatSetValuesCOO_MPIAIJHIPSPARSE()
207   if (coo->sendlen) {  in MatSetValuesCOO_MPIAIJHIPSPARSE()
208 …IP_KERNEL_NAME(MatPackCOOValues), dim3((coo->sendlen + 255) / 256), dim3(256), 0, PetscDefaultHipS…  in MatSetValuesCOO_MPIAIJHIPSPARSE()
213 …PetscCall(PetscSFReduceWithMemTypeBegin(coo->sf, MPIU_SCALAR, PETSC_MEMTYPE_HIP, vsend, PETSC_MEMT…  in MatSetValuesCOO_MPIAIJHIPSPARSE()
219   PetscCall(PetscSFReduceEnd(coo->sf, MPIU_SCALAR, vsend, v2, MPI_REPLACE));  in MatSetValuesCOO_MPIAIJHIPSPARSE()