xref: /petsc/src/mat/impls/hypre/mhypre_kernels.hpp (revision b5b1a1b87f6fcd196fc1e33da00be2a37963df95)
1 
2 #include <../src/mat/impls/hypre/mhypre.h>
3 
4 // Zero the specified n rows in rows[] of the hypre CSRMatrix (i, j, a) and replace the diagonal entry with diag
5 __global__ static void ZeroRows(PetscInt n, const PetscInt rows[], const HYPRE_Int i[], const HYPRE_Int j[], HYPRE_Complex a[], HYPRE_Complex diag)
6 {
7   PetscInt k     = blockDim.x * blockIdx.x + threadIdx.x; // k-th entry in rows[]
8   PetscInt c     = blockDim.y * blockIdx.y + threadIdx.y; // c-th nonzero in row rows[k]
9   PetscInt gridx = gridDim.x * blockDim.x;
10   PetscInt gridy = gridDim.y * blockDim.y;
11   for (; k < n; k += gridx) {
12     PetscInt r  = rows[k]; // r-th row of the matrix
13     PetscInt nz = i[r + 1] - i[r];
14     for (; c < nz; c += gridy) {
15       if (r == j[i[r] + c]) a[i[r] + c] = diag;
16       else a[i[r] + c] = 0.0;
17     }
18   }
19 }
20