1735d7f90SBarry Smith #include <petscdevice.h> 2735d7f90SBarry Smith #include "ex18.h" 3735d7f90SBarry Smith 4*9371c9d4SSatish Balay __global__ void FillValues(PetscInt n, PetscScalar *v) { 5735d7f90SBarry Smith PetscInt i = blockIdx.x * blockDim.x + threadIdx.x; 6735d7f90SBarry Smith PetscScalar *s; 7735d7f90SBarry Smith if (i < n) { 8735d7f90SBarry Smith s = &v[3 * 3 * i]; 9735d7f90SBarry Smith for (PetscInt vi = 0; vi < 3; vi++) { 10*9371c9d4SSatish Balay for (PetscInt vj = 0; vj < 3; vj++) { s[vi * 3 + vj] = vi + 2 * vj; } 11735d7f90SBarry Smith } 12735d7f90SBarry Smith } 13735d7f90SBarry Smith } 14735d7f90SBarry Smith 15*9371c9d4SSatish Balay PetscErrorCode FillMatrixCUDACOO(FEStruct *fe, Mat A) { 16735d7f90SBarry Smith PetscScalar *v; 17735d7f90SBarry Smith 18735d7f90SBarry Smith PetscFunctionBeginUser; 199566063dSJacob Faibussowitsch PetscCallCUDA(cudaMalloc((void **)&v, 3 * 3 * fe->Ne * sizeof(PetscScalar))); 20735d7f90SBarry Smith FillValues<<<(fe->Ne + 255) / 256, 256>>>(fe->Ne, v); 219566063dSJacob Faibussowitsch PetscCall(MatSetValuesCOO(A, v, INSERT_VALUES)); 229566063dSJacob Faibussowitsch PetscCallCUDA(cudaFree(v)); 23735d7f90SBarry Smith PetscFunctionReturn(0); 24735d7f90SBarry Smith } 25