44 static PetscErrorCode MatSetPreallocationCOO_MPIAIJCUSPARSE(Mat mat, PetscCount coo_n, PetscInt coo…  in MatSetPreallocationCOO_MPIAIJCUSPARSE()
89   PetscCallCUDA(cudaMalloc((void **)&coo_d->Ajmap1, (coo_h->Annz + 1) * sizeof(PetscCount)));  in MatSetPreallocationCOO_MPIAIJCUSPARSE()
90   PetscCallCUDA(cudaMalloc((void **)&coo_d->Aperm1, coo_h->Atot1 * sizeof(PetscCount)));  in MatSetPreallocationCOO_MPIAIJCUSPARSE()
91   PetscCallCUDA(cudaMalloc((void **)&coo_d->Bjmap1, (coo_h->Bnnz + 1) * sizeof(PetscCount)));  in MatSetPreallocationCOO_MPIAIJCUSPARSE()
92   PetscCallCUDA(cudaMalloc((void **)&coo_d->Bperm1, coo_h->Btot1 * sizeof(PetscCount)));  in MatSetPreallocationCOO_MPIAIJCUSPARSE()
93   PetscCallCUDA(cudaMalloc((void **)&coo_d->Aimap2, coo_h->Annz2 * sizeof(PetscCount)));  in MatSetPreallocationCOO_MPIAIJCUSPARSE()
94   PetscCallCUDA(cudaMalloc((void **)&coo_d->Ajmap2, (coo_h->Annz2 + 1) * sizeof(PetscCount)));  in MatSetPreallocationCOO_MPIAIJCUSPARSE()
95   PetscCallCUDA(cudaMalloc((void **)&coo_d->Aperm2, coo_h->Atot2 * sizeof(PetscCount)));  in MatSetPreallocationCOO_MPIAIJCUSPARSE()
96   PetscCallCUDA(cudaMalloc((void **)&coo_d->Bimap2, coo_h->Bnnz2 * sizeof(PetscCount)));  in MatSetPreallocationCOO_MPIAIJCUSPARSE()
97   PetscCallCUDA(cudaMalloc((void **)&coo_d->Bjmap2, (coo_h->Bnnz2 + 1) * sizeof(PetscCount)));  in MatSetPreallocationCOO_MPIAIJCUSPARSE()
98   PetscCallCUDA(cudaMalloc((void **)&coo_d->Bperm2, coo_h->Btot2 * sizeof(PetscCount)));  in MatSetPreallocationCOO_MPIAIJCUSPARSE()
99   PetscCallCUDA(cudaMalloc((void **)&coo_d->Cperm1, coo_h->sendlen * sizeof(PetscCount)));  in MatSetPreallocationCOO_MPIAIJCUSPARSE()
103 …PetscCallCUDA(cudaMemcpy(coo_d->Ajmap1, coo_h->Ajmap1, (coo_h->Annz + 1) * sizeof(PetscCount), cud…  in MatSetPreallocationCOO_MPIAIJCUSPARSE()
104 …PetscCallCUDA(cudaMemcpy(coo_d->Aperm1, coo_h->Aperm1, coo_h->Atot1 * sizeof(PetscCount), cudaMemc…  in MatSetPreallocationCOO_MPIAIJCUSPARSE()
105 …PetscCallCUDA(cudaMemcpy(coo_d->Bjmap1, coo_h->Bjmap1, (coo_h->Bnnz + 1) * sizeof(PetscCount), cud…  in MatSetPreallocationCOO_MPIAIJCUSPARSE()
106 …PetscCallCUDA(cudaMemcpy(coo_d->Bperm1, coo_h->Bperm1, coo_h->Btot1 * sizeof(PetscCount), cudaMemc…  in MatSetPreallocationCOO_MPIAIJCUSPARSE()
107 …PetscCallCUDA(cudaMemcpy(coo_d->Aimap2, coo_h->Aimap2, coo_h->Annz2 * sizeof(PetscCount), cudaMemc…  in MatSetPreallocationCOO_MPIAIJCUSPARSE()
108 …PetscCallCUDA(cudaMemcpy(coo_d->Ajmap2, coo_h->Ajmap2, (coo_h->Annz2 + 1) * sizeof(PetscCount), cu…  in MatSetPreallocationCOO_MPIAIJCUSPARSE()
109 …PetscCallCUDA(cudaMemcpy(coo_d->Aperm2, coo_h->Aperm2, coo_h->Atot2 * sizeof(PetscCount), cudaMemc…  in MatSetPreallocationCOO_MPIAIJCUSPARSE()
110 …PetscCallCUDA(cudaMemcpy(coo_d->Bimap2, coo_h->Bimap2, coo_h->Bnnz2 * sizeof(PetscCount), cudaMemc…  in MatSetPreallocationCOO_MPIAIJCUSPARSE()
111 …PetscCallCUDA(cudaMemcpy(coo_d->Bjmap2, coo_h->Bjmap2, (coo_h->Bnnz2 + 1) * sizeof(PetscCount), cu…  in MatSetPreallocationCOO_MPIAIJCUSPARSE()
112 …PetscCallCUDA(cudaMemcpy(coo_d->Bperm2, coo_h->Bperm2, coo_h->Btot2 * sizeof(PetscCount), cudaMemc…  in MatSetPreallocationCOO_MPIAIJCUSPARSE()
113 …PetscCallCUDA(cudaMemcpy(coo_d->Cperm1, coo_h->Cperm1, coo_h->sendlen * sizeof(PetscCount), cudaMe…  in MatSetPreallocationCOO_MPIAIJCUSPARSE()
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()