1: static char help[] = "Solves a tridiagonal linear system with CUDA managed memory.\n\n";

  3: #include <petscdevice_cuda.h>
  4: #include <petscksp.h>

  6: // adapted from ksp/tutorials/ex23.c

  8: int main(int argc, char **args)
  9: {
 10:   Vec          x, b, u;
 11:   Mat          A;
 12:   KSP          ksp;
 13:   PetscReal    norm, tol = 1000. * PETSC_MACHINE_EPSILON;
 14:   PetscInt     i, n, N = 32, col[3], its, rstart, rend;
 15:   PetscScalar  value[3];
 16:   PetscScalar *xarray, *uarray;
 17:   PetscScalar *array;

 19:   PetscFunctionBeginUser;
 20:   PetscCall(PetscInitialize(&argc, &args, NULL, help));
 21:   PetscCall(PetscOptionsGetInt(NULL, NULL, "-N", &N, NULL));
 22:   PetscCall(MatCreate(PETSC_COMM_WORLD, &A));
 23:   PetscCall(MatSetSizes(A, PETSC_DECIDE, PETSC_DECIDE, N, N));
 24:   PetscCall(MatSetType(A, MATAIJCUSPARSE));
 25:   PetscCall(MatSetFromOptions(A));
 26:   PetscCall(MatSetUp(A));
 27:   PetscCall(MatGetLocalSize(A, &n, NULL));
 28:   PetscCall(MatGetOwnershipRange(A, &rstart, &rend)); // Same row/column range for this test

 30:   if (!rstart) {
 31:     rstart   = 1;
 32:     i        = 0;
 33:     col[0]   = 0;
 34:     col[1]   = 1;
 35:     value[0] = 2.0;
 36:     value[1] = -1.0;
 37:     PetscCall(MatSetValues(A, 1, &i, 2, col, value, INSERT_VALUES));
 38:   }
 39:   if (rend == N) {
 40:     rend     = N - 1;
 41:     i        = N - 1;
 42:     col[0]   = N - 2;
 43:     col[1]   = N - 1;
 44:     value[0] = -1.0;
 45:     value[1] = 2.0;
 46:     PetscCall(MatSetValues(A, 1, &i, 2, col, value, INSERT_VALUES));
 47:   }

 49:   value[0] = -1.0;
 50:   value[1] = 2.0;
 51:   value[2] = -1.0;
 52:   for (i = rstart; i < rend; i++) {
 53:     col[0] = i - 1;
 54:     col[1] = i;
 55:     col[2] = i + 1;
 56:     PetscCall(MatSetValues(A, 1, &i, 3, col, value, INSERT_VALUES));
 57:   }

 59:   PetscCall(MatAssemblyBegin(A, MAT_FINAL_ASSEMBLY));
 60:   PetscCall(MatAssemblyEnd(A, MAT_FINAL_ASSEMBLY));

 62:   PetscCallCUDA(cudaMallocManaged((void **)&xarray, n * sizeof(PetscScalar), cudaMemAttachGlobal));
 63:   PetscCallCUDA(cudaMallocManaged((void **)&uarray, n * sizeof(PetscScalar), cudaMemAttachGlobal));

 65:   // Use a managed array as a host array parameter, it will be a host array in petsc's view;
 66:   // similarily, use it as a device array parameter, it will be a device array.
 67:   // One can use the same array on both host and device parameters.
 68:   PetscCall(VecCreateMPICUDAWithArrays(PETSC_COMM_WORLD, 1, n, N, xarray, xarray, &x));
 69:   PetscCall(VecCreateMPICUDAWithArrays(PETSC_COMM_WORLD, 1, n, N, uarray, uarray, &u));

 71:   PetscCall(MatCreateVecs(A, NULL, &b));
 72:   PetscCall(VecSet(u, 1.0)); // Set u on device

 74:   PetscCall(VecGetArray(u, &array)); // Do a cudaMemcpyDeviceToHost from uarray to uarray!
 75:   PetscCheck(array == uarray && array[0] == 1.0, PETSC_COMM_SELF, PETSC_ERR_PLIB, "u is out of sync");
 76:   array[0] = 2.0;
 77:   PetscCall(VecRestoreArray(u, &array));
 78:   PetscCall(MatMult(A, u, b));

 80:   PetscCall(KSPCreate(PETSC_COMM_WORLD, &ksp));
 81:   PetscCall(KSPSetOperators(ksp, A, A));
 82:   PetscCall(KSPSetFromOptions(ksp));
 83:   PetscCall(KSPSolve(ksp, b, x));

 85:   PetscCall(VecAXPY(x, -1.0, u));
 86:   PetscCall(VecNorm(x, NORM_2, &norm));
 87:   PetscCall(KSPGetIterationNumber(ksp, &its));
 88:   if (norm > tol) PetscCall(PetscPrintf(PETSC_COMM_WORLD, "Norm of error %g, Iterations %" PetscInt_FMT "\n", (double)norm, its));

 90:   PetscCallCUDA(cudaFree(xarray));
 91:   PetscCallCUDA(cudaFree(uarray));
 92:   PetscCall(VecDestroy(&x));
 93:   PetscCall(VecDestroy(&u));
 94:   PetscCall(VecDestroy(&b));
 95:   PetscCall(MatDestroy(&A));
 96:   PetscCall(KSPDestroy(&ksp));
 97:   PetscCall(PetscFinalize());
 98:   return 0;
 99: }

101: /*TEST

103:   build:
104:     requires: cuda !complex !single

106:   test:
107:     nsize: {{1 2}}
108:     output_file: output/empty.out

110: TEST*/