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