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