1*78923ab9SStefano Zampini #include <petsc/private/petschypre.h> 2a32e9c99SJunchao Zhang #include <petscdevice_cuda.h> 3a32e9c99SJunchao Zhang #include <../src/mat/impls/hypre/mhypre_kernels.hpp> 4a32e9c99SJunchao Zhang 5a32e9c99SJunchao Zhang PetscErrorCode MatZeroRows_CUDA(PetscInt n, const PetscInt rows[], const HYPRE_Int i[], const HYPRE_Int j[], HYPRE_Complex a[], HYPRE_Complex diag) 6a32e9c99SJunchao Zhang { 7a32e9c99SJunchao Zhang const PetscInt blkDimX = 16, blkDimY = 32; 8a32e9c99SJunchao Zhang PetscInt gridDimX = (n + blkDimX - 1) / blkDimX; 9a32e9c99SJunchao Zhang cudaStream_t stream; 10a32e9c99SJunchao Zhang 11a32e9c99SJunchao Zhang PetscFunctionBegin; 12a32e9c99SJunchao Zhang if (!n) PetscFunctionReturn(PETSC_SUCCESS); 13a32e9c99SJunchao Zhang PetscCall(PetscGetCurrentCUDAStream(&stream)); 14a32e9c99SJunchao Zhang ZeroRows<<<dim3(gridDimX, 1), dim3(blkDimX, blkDimY), 0, stream>>>(n, rows, i, j, a, diag); 15a32e9c99SJunchao Zhang PetscCallCUDA(cudaGetLastError()); 16a32e9c99SJunchao Zhang PetscFunctionReturn(PETSC_SUCCESS); 17a32e9c99SJunchao Zhang } 18