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