44 static PetscErrorCode MatSetPreallocationCOO_MPIAIJHIPSPARSE(Mat mat, PetscCount coo_n, PetscInt co…  in MatSetPreallocationCOO_MPIAIJHIPSPARSE()
89   PetscCallHIP(hipMalloc((void **)&coo_d->Ajmap1, (coo_h->Annz + 1) * sizeof(PetscCount)));  in MatSetPreallocationCOO_MPIAIJHIPSPARSE()
90   PetscCallHIP(hipMalloc((void **)&coo_d->Aperm1, coo_h->Atot1 * sizeof(PetscCount)));  in MatSetPreallocationCOO_MPIAIJHIPSPARSE()
91   PetscCallHIP(hipMalloc((void **)&coo_d->Bjmap1, (coo_h->Bnnz + 1) * sizeof(PetscCount)));  in MatSetPreallocationCOO_MPIAIJHIPSPARSE()
92   PetscCallHIP(hipMalloc((void **)&coo_d->Bperm1, coo_h->Btot1 * sizeof(PetscCount)));  in MatSetPreallocationCOO_MPIAIJHIPSPARSE()
93   PetscCallHIP(hipMalloc((void **)&coo_d->Aimap2, coo_h->Annz2 * sizeof(PetscCount)));  in MatSetPreallocationCOO_MPIAIJHIPSPARSE()
94   PetscCallHIP(hipMalloc((void **)&coo_d->Ajmap2, (coo_h->Annz2 + 1) * sizeof(PetscCount)));  in MatSetPreallocationCOO_MPIAIJHIPSPARSE()
95   PetscCallHIP(hipMalloc((void **)&coo_d->Aperm2, coo_h->Atot2 * sizeof(PetscCount)));  in MatSetPreallocationCOO_MPIAIJHIPSPARSE()
96   PetscCallHIP(hipMalloc((void **)&coo_d->Bimap2, coo_h->Bnnz2 * sizeof(PetscCount)));  in MatSetPreallocationCOO_MPIAIJHIPSPARSE()
97   PetscCallHIP(hipMalloc((void **)&coo_d->Bjmap2, (coo_h->Bnnz2 + 1) * sizeof(PetscCount)));  in MatSetPreallocationCOO_MPIAIJHIPSPARSE()
98   PetscCallHIP(hipMalloc((void **)&coo_d->Bperm2, coo_h->Btot2 * sizeof(PetscCount)));  in MatSetPreallocationCOO_MPIAIJHIPSPARSE()
99   PetscCallHIP(hipMalloc((void **)&coo_d->Cperm1, coo_h->sendlen * sizeof(PetscCount)));  in MatSetPreallocationCOO_MPIAIJHIPSPARSE()
103 …PetscCallHIP(hipMemcpy(coo_d->Ajmap1, coo_h->Ajmap1, (coo_h->Annz + 1) * sizeof(PetscCount), hipMe…  in MatSetPreallocationCOO_MPIAIJHIPSPARSE()
104 …PetscCallHIP(hipMemcpy(coo_d->Aperm1, coo_h->Aperm1, coo_h->Atot1 * sizeof(PetscCount), hipMemcpyH…  in MatSetPreallocationCOO_MPIAIJHIPSPARSE()
105 …PetscCallHIP(hipMemcpy(coo_d->Bjmap1, coo_h->Bjmap1, (coo_h->Bnnz + 1) * sizeof(PetscCount), hipMe…  in MatSetPreallocationCOO_MPIAIJHIPSPARSE()
106 …PetscCallHIP(hipMemcpy(coo_d->Bperm1, coo_h->Bperm1, coo_h->Btot1 * sizeof(PetscCount), hipMemcpyH…  in MatSetPreallocationCOO_MPIAIJHIPSPARSE()
107 …PetscCallHIP(hipMemcpy(coo_d->Aimap2, coo_h->Aimap2, coo_h->Annz2 * sizeof(PetscCount), hipMemcpyH…  in MatSetPreallocationCOO_MPIAIJHIPSPARSE()
108 …PetscCallHIP(hipMemcpy(coo_d->Ajmap2, coo_h->Ajmap2, (coo_h->Annz2 + 1) * sizeof(PetscCount), hipM…  in MatSetPreallocationCOO_MPIAIJHIPSPARSE()
109 …PetscCallHIP(hipMemcpy(coo_d->Aperm2, coo_h->Aperm2, coo_h->Atot2 * sizeof(PetscCount), hipMemcpyH…  in MatSetPreallocationCOO_MPIAIJHIPSPARSE()
110 …PetscCallHIP(hipMemcpy(coo_d->Bimap2, coo_h->Bimap2, coo_h->Bnnz2 * sizeof(PetscCount), hipMemcpyH…  in MatSetPreallocationCOO_MPIAIJHIPSPARSE()
111 …PetscCallHIP(hipMemcpy(coo_d->Bjmap2, coo_h->Bjmap2, (coo_h->Bnnz2 + 1) * sizeof(PetscCount), hipM…  in MatSetPreallocationCOO_MPIAIJHIPSPARSE()
112 …PetscCallHIP(hipMemcpy(coo_d->Bperm2, coo_h->Bperm2, coo_h->Btot2 * sizeof(PetscCount), hipMemcpyH…  in MatSetPreallocationCOO_MPIAIJHIPSPARSE()
113 …PetscCallHIP(hipMemcpy(coo_d->Cperm1, coo_h->Cperm1, coo_h->sendlen * sizeof(PetscCount), hipMemcp…  in MatSetPreallocationCOO_MPIAIJHIPSPARSE()
120 __global__ static void MatPackCOOValues(const PetscScalar kv[], PetscCount nnz, const PetscCount pe…  in MatPackCOOValues()
122   PetscCount       i         = blockIdx.x * blockDim.x + threadIdx.x;  in MatPackCOOValues()
123   const PetscCount grid_size = gridDim.x * blockDim.x;  in MatPackCOOValues()
127 …, PetscCount Annz, const PetscCount Ajmap1[], const PetscCount Aperm1[], PetscScalar Aa[], PetscCo…  in MatAddLocalCOOValues()
129   PetscCount       i         = blockIdx.x * blockDim.x + threadIdx.x;  in MatAddLocalCOOValues()
130   const PetscCount grid_size = gridDim.x * blockDim.x;  in MatAddLocalCOOValues()
134       for (PetscCount k = Ajmap1[i]; k < Ajmap1[i + 1]; k++) sum += kv[Aperm1[k]];  in MatAddLocalCOOValues()
138       for (PetscCount k = Bjmap1[i]; k < Bjmap1[i + 1]; k++) sum += kv[Bperm1[k]];  in MatAddLocalCOOValues()
144 …PetscCount Annz2, const PetscCount Aimap2[], const PetscCount Ajmap2[], const PetscCount Aperm2[],…  in MatAddRemoteCOOValues()
146   PetscCount       i         = blockIdx.x * blockDim.x + threadIdx.x;  in MatAddRemoteCOOValues()
147   const PetscCount grid_size = gridDim.x * blockDim.x;  in MatAddRemoteCOOValues()
150       for (PetscCount k = Ajmap2[i]; k < Ajmap2[i + 1]; k++) Aa[Aimap2[i]] += kv[Aperm2[k]];  in MatAddRemoteCOOValues()
153       for (PetscCount k = Bjmap2[i]; k < Bjmap2[i + 1]; k++) Ba[Bimap2[i]] += kv[Bperm2[k]];  in MatAddRemoteCOOValues()