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
MatZeroRows_HIP(PetscInt n,const PetscInt rows[],const HYPRE_Int i[],const HYPRE_Int j[],HYPRE_Complex a[],HYPRE_Complex diag)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