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