Merge branch 'barry/2025-10-18/optimize-matzeroentries-seqaijcusparse' into 'main'

MatZeroEntries_SeqAIJCUSPARSE() was zeroing both the CPU and GPU array....

See merge request petsc/petsc!8796

Author: balay

src/mat/impls/aij/seq/seqcusparse/aijcusparse.cu

@@ -3924,16 +3924,16 @@ static PetscErrorCode MatScale_SeqAIJCUSPARSE(Mat Y, PetscScalar a)
 
 static PetscErrorCode MatZeroEntries_SeqAIJCUSPARSE(Mat A)
 {
-  PetscBool   both = PETSC_FALSE;
-  Mat_SeqAIJ *a    = (Mat_SeqAIJ *)A->data;
+  PetscBool   gpu = PETSC_FALSE;
+  Mat_SeqAIJ *a   = (Mat_SeqAIJ *)A->data;
 
   PetscFunctionBegin;
   if (A->factortype == MAT_FACTOR_NONE) {
     Mat_SeqAIJCUSPARSE *spptr = (Mat_SeqAIJCUSPARSE *)A->spptr;
     if (spptr->mat) {
       CsrMatrix *matrix = (CsrMatrix *)spptr->mat->mat;
       if (matrix->values) {
-        both = PETSC_TRUE;
+        gpu = PETSC_TRUE;
         thrust::fill(thrust::device, matrix->values->begin(), matrix->values->end(), 0.);
       }
     }
@@ -3942,10 +3942,12 @@ static PetscErrorCode MatZeroEntries_SeqAIJCUSPARSE(Mat A)
       if (matrix->values) thrust::fill(thrust::device, matrix->values->begin(), matrix->values->end(), 0.);
     }
   }
-  PetscCall(PetscArrayzero(a->a, a->i[A->rmap->n]));
+  if (gpu) A->offloadmask = PETSC_OFFLOAD_GPU;
+  else {
+    PetscCall(PetscArrayzero(a->a, a->i[A->rmap->n]));
+    A->offloadmask = PETSC_OFFLOAD_CPU;
+  }
   PetscCall(MatSeqAIJInvalidateDiagonal(A));
-  if (both) A->offloadmask = PETSC_OFFLOAD_BOTH;
-  else A->offloadmask = PETSC_OFFLOAD_CPU;
   PetscFunctionReturn(PETSC_SUCCESS);
 }