1 #include <cuda_runtime.h> 2 #include <petscdevice.h> 3 #include "ex18.h" 4 5 __global__ void FillValues(PetscInt n, PetscScalar *v) 6 { 7 PetscInt i = blockIdx.x * blockDim.x + threadIdx.x; 8 PetscScalar *s; 9 if (i < n) { 10 s = &v[3*3*i]; 11 for (PetscInt vi=0; vi<3; vi++) { 12 for (PetscInt vj=0; vj<3; vj++) { 13 s[vi*3+vj] = vi+2*vj; 14 } 15 } 16 } 17 } 18 19 PetscErrorCode FillMatrixCUDACOO(FEStruct *fe,Mat A) 20 { 21 PetscErrorCode ierr; 22 cudaError_t cerr; 23 PetscScalar *v; 24 25 PetscFunctionBeginUser; 26 cerr = cudaMalloc((void**)&v,3*3*fe->Ne*sizeof(PetscScalar));CHKERRCUDA(cerr); 27 FillValues<<<(fe->Ne+255)/256,256>>>(fe->Ne,v); 28 ierr = MatSetValuesCOO(A,v,INSERT_VALUES);CHKERRQ(ierr); 29 cerr = cudaFree(v);CHKERRCUDA(cerr); 30 PetscFunctionReturn(0); 31 } 32