GCC Code Coverage Report

Directory:	./
File:	src/svd/impls/cyclic/cycliccuda/cycliccuda.cu
Date:	2026-10-22 02:08:06

	Exec	Total	Coverage
Lines:	79	79	100.0%
Functions:	3	4	75.0%
Branches:	81	140	57.9%

  
      Line
      Branch
      Exec
      Source
    
      /*
    
         - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - -
    
         SLEPc - Scalable Library for Eigenvalue Problem Computations
    
         Copyright (c) 2002-, Universitat Politecnica de Valencia, Spain
    
         This file is part of SLEPc.
    
         SLEPc is distributed under a 2-clause BSD license (see LICENSE).
    
         - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - -
    
      */
    
      /*
    
         SLEPc singular value solver: "cyclic" (CUDA implementation)
    
      */
    
      #include <slepc/private/svdimpl.h>
    
      #include <petscdevice_cuda.h>
    
      #include "../src/svd/impls/cyclic/cyclic.h"
    
      /* block alignment can differ between ranks, so copies must not log collective VecCopy() events */
    
      5162
      static PetscErrorCode Copy_Cyclic_CUDA(PetscScalar *dest,const PetscScalar *src,PetscInt n)
    
      {
    
      5162
        cudaStream_t stream;
    
      5162
        PetscFunctionBegin;
    
        1/2✓ Branch 0 taken 1 times.
✗ Branch 1 not taken.

      5162
        if (n) {
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 1 times.

      5162
          PetscCall(PetscGetCurrentCUDAStream(&stream));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 1 times.

      5162
          PetscCall(PetscLogGpuTimeBegin());
    
        1/4✗ Branch 0 not taken.
✓ Branch 1 taken 1 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.

      5162
          PetscCallCUDA(cudaMemcpyAsync(dest,src,n*sizeof(*src),cudaMemcpyDeviceToDevice,stream));
    
        1/2✓ Branch 0 taken 1 times.
✗ Branch 1 not taken.

      5162
          PetscCall(PetscLogGpuTimeEnd());
    
        }
    
        PetscFunctionReturn(PETSC_SUCCESS);
    
      }
    
      348
      PetscErrorCode MatMult_Cyclic_CUDA(Mat B,Vec x,Vec y)
    
      {
    
      348
        SVD_CYCLIC_SHELL  *ctx;
    
      348
        const PetscScalar *d_px,*read;
    
      348
        PetscScalar       *d_py,*write;
    
      348
        PetscInt          m,n;
    
      348
        PetscFunctionBegin;
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.

      348
        PetscCall(MatShellGetContext(B,&ctx));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.

      348
        PetscCall(MatGetLocalSize(ctx->A,&m,&n));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.

      348
        PetscCall(VecCUDAGetArrayRead(x,&d_px));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.

      348
        PetscCall(VecCUDAGetArrayWrite(y,&d_py));
    
        /* dense matrix columns may misalign either block; retain the same Vec objects on every rank */
    
        2/2✓ Branch 0 taken 1 times.
✓ Branch 1 taken 2 times.

      348
        if (((PETSC_UINTPTR_T)(d_px))%16) {
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 1 times.

      6
          PetscCall(VecCUDAGetArrayWrite(ctx->x1,&write));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 1 times.

      6
          PetscCall(Copy_Cyclic_CUDA(write,d_px,m));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 1 times.

      6
          PetscCall(VecCUDARestoreArrayWrite(ctx->x1,&write));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.

      342
        } else PetscCall(VecCUDAPlaceArray(ctx->x1,d_px));
    
        2/2✓ Branch 0 taken 1 times.
✓ Branch 1 taken 2 times.

      348
        if (((PETSC_UINTPTR_T)(d_px+m))%16) {
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 1 times.

      167
          PetscCall(VecCUDAGetArrayWrite(ctx->x2,&write));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 1 times.

      167
          PetscCall(Copy_Cyclic_CUDA(write,d_px+m,n));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 1 times.

      167
          PetscCall(VecCUDARestoreArrayWrite(ctx->x2,&write));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.

      181
        } else PetscCall(VecCUDAPlaceArray(ctx->x2,d_px+m));
    
        3/4✓ Branch 0 taken 2 times.
✓ Branch 1 taken 1 times.
✗ Branch 2 not taken.
✓ Branch 3 taken 2 times.

      348
        if (!(((PETSC_UINTPTR_T)(d_py))%16)) PetscCall(VecCUDAPlaceArray(ctx->y1,d_py));
    
        3/4✓ Branch 0 taken 2 times.
✓ Branch 1 taken 1 times.
✗ Branch 2 not taken.
✓ Branch 3 taken 2 times.

      348
        if (!(((PETSC_UINTPTR_T)(d_py+m))%16)) PetscCall(VecCUDAPlaceArray(ctx->y2,d_py+m));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.

      348
        PetscCall(MatMult(ctx->A,ctx->x2,ctx->y1));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.

      348
        PetscCall(MatMult(ctx->AT,ctx->x1,ctx->y2));
    
        2/2✓ Branch 0 taken 1 times.
✓ Branch 1 taken 2 times.

      348
        if (((PETSC_UINTPTR_T)(d_py))%16) {
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 1 times.

      6
          PetscCall(VecCUDAGetArrayRead(ctx->y1,&read));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 1 times.

      6
          PetscCall(Copy_Cyclic_CUDA(d_py,read,m));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 1 times.

      6
          PetscCall(VecCUDARestoreArrayRead(ctx->y1,&read));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.

      342
        } else PetscCall(VecCUDAResetArray(ctx->y1));
    
        2/2✓ Branch 0 taken 1 times.
✓ Branch 1 taken 2 times.

      348
        if (((PETSC_UINTPTR_T)(d_py+m))%16) {
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 1 times.

      167
          PetscCall(VecCUDAGetArrayRead(ctx->y2,&read));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 1 times.

      167
          PetscCall(Copy_Cyclic_CUDA(d_py+m,read,n));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 1 times.

      167
          PetscCall(VecCUDARestoreArrayRead(ctx->y2,&read));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.

      181
        } else PetscCall(VecCUDAResetArray(ctx->y2));
    
        3/4✓ Branch 0 taken 2 times.
✓ Branch 1 taken 1 times.
✗ Branch 2 not taken.
✓ Branch 3 taken 2 times.

      348
        if (!(((PETSC_UINTPTR_T)(d_px))%16)) PetscCall(VecCUDAResetArray(ctx->x1));
    
        3/4✓ Branch 0 taken 2 times.
✓ Branch 1 taken 1 times.
✗ Branch 2 not taken.
✓ Branch 3 taken 2 times.

      348
        if (!(((PETSC_UINTPTR_T)(d_px+m))%16)) PetscCall(VecCUDAResetArray(ctx->x2));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.

      348
        PetscCall(VecCUDARestoreArrayRead(x,&d_px));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.

      348
        PetscCall(VecCUDARestoreArrayWrite(y,&d_py));
    
        PetscFunctionReturn(PETSC_SUCCESS);
    
      }
    
      4861
      PetscErrorCode MatMult_ECross_CUDA(Mat B,Vec x,Vec y)
    
      {
    
      4861
        SVD_CYCLIC_SHELL  *ctx;
    
      4861
        const PetscScalar *d_px,*read;
    
      4861
        PetscScalar       *d_py,*write;
    
      4861
        PetscInt          mn,m,n;
    
      4861
        PetscFunctionBegin;
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.

      4861
        PetscCall(MatShellGetContext(B,&ctx));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.

      4861
        PetscCall(MatGetLocalSize(ctx->A,NULL,&n));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.

      4861
        PetscCall(VecGetLocalSize(y,&mn));
    
      4861
        m = mn-n;
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.

      4861
        PetscCall(VecCUDAGetArrayRead(x,&d_px));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.

      4861
        PetscCall(VecCUDAGetArrayWrite(y,&d_py));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.

      4861
        PetscCall(VecCUDAPlaceArray(ctx->x1,d_px));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.

      4861
        PetscCall(VecCUDAPlaceArray(ctx->y1,d_py));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.

      4861
        PetscCall(VecCopy(ctx->x1,ctx->y1));
    
        2/2✓ Branch 0 taken 1 times.
✓ Branch 1 taken 1 times.

      4861
        if (((PETSC_UINTPTR_T)(d_px+m))%16) {
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 1 times.

      2408
          PetscCall(VecCUDAGetArrayWrite(ctx->x2,&write));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 1 times.

      2408
          PetscCall(Copy_Cyclic_CUDA(write,d_px+m,n));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 1 times.

      2408
          PetscCall(VecCUDARestoreArrayWrite(ctx->x2,&write));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 1 times.

      2453
        } else PetscCall(VecCUDAPlaceArray(ctx->x2,d_px+m));
    
        3/4✓ Branch 0 taken 1 times.
✓ Branch 1 taken 1 times.
✗ Branch 2 not taken.
✓ Branch 3 taken 1 times.

      4861
        if (!(((PETSC_UINTPTR_T)(d_py+m))%16)) PetscCall(VecCUDAPlaceArray(ctx->y2,d_py+m));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.

      4861
        PetscCall(MatMult(ctx->A,ctx->x2,ctx->w));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.

      4861
        PetscCall(MatMult(ctx->AT,ctx->w,ctx->y2));
    
        2/2✓ Branch 0 taken 1 times.
✓ Branch 1 taken 1 times.

      4861
        if (((PETSC_UINTPTR_T)(d_py+m))%16) {
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 1 times.

      2408
          PetscCall(VecCUDAGetArrayRead(ctx->y2,&read));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 1 times.

      2408
          PetscCall(Copy_Cyclic_CUDA(d_py+m,read,n));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 1 times.

      2408
          PetscCall(VecCUDARestoreArrayRead(ctx->y2,&read));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 1 times.

      2453
        } else PetscCall(VecCUDAResetArray(ctx->y2));
    
        3/4✓ Branch 0 taken 1 times.
✓ Branch 1 taken 1 times.
✗ Branch 2 not taken.
✓ Branch 3 taken 1 times.

      4861
        if (!(((PETSC_UINTPTR_T)(d_px+m))%16)) PetscCall(VecCUDAResetArray(ctx->x2));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.

      4861
        PetscCall(VecCUDAResetArray(ctx->x1));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.

      4861
        PetscCall(VecCUDAResetArray(ctx->y1));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.

      4861
        PetscCall(VecCUDARestoreArrayRead(x,&d_px));
    
        1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.

      4861
        PetscCall(VecCUDARestoreArrayWrite(y,&d_py));
    
        PetscFunctionReturn(PETSC_SUCCESS);
    
      }

Line	Branch	Exec	Source
1			/*
2			- - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - -
3			SLEPc - Scalable Library for Eigenvalue Problem Computations
4			Copyright (c) 2002-, Universitat Politecnica de Valencia, Spain
5
6			This file is part of SLEPc.
7			SLEPc is distributed under a 2-clause BSD license (see LICENSE).
8			- - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - -
9			*/
10			/*
11			SLEPc singular value solver: "cyclic" (CUDA implementation)
12			*/
13			#include <slepc/private/svdimpl.h>
14			#include <petscdevice_cuda.h>
15			#include "../src/svd/impls/cyclic/cyclic.h"
16
17			/* block alignment can differ between ranks, so copies must not log collective VecCopy() events */
18		5162	static PetscErrorCode Copy_Cyclic_CUDA(PetscScalar dest,const PetscScalar src,PetscInt n)
19			{
20		5162	cudaStream_t stream;
21
22		5162	PetscFunctionBegin;
23	1/2 ✓ Branch 0 taken 1 times. ✗ Branch 1 not taken.	5162	if (n) {
24	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 1 times.	5162	PetscCall(PetscGetCurrentCUDAStream(&stream));
25	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 1 times.	5162	PetscCall(PetscLogGpuTimeBegin());
26	1/4 ✗ Branch 0 not taken. ✓ Branch 1 taken 1 times. ✗ Branch 2 not taken. ✗ Branch 3 not taken.	5162	PetscCallCUDA(cudaMemcpyAsync(dest,src,nsizeof(src),cudaMemcpyDeviceToDevice,stream));
27	1/2 ✓ Branch 0 taken 1 times. ✗ Branch 1 not taken.	5162	PetscCall(PetscLogGpuTimeEnd());
28			}
29			PetscFunctionReturn(PETSC_SUCCESS);
30			}
31
32		348	PetscErrorCode MatMult_Cyclic_CUDA(Mat B,Vec x,Vec y)
33			{
34		348	SVD_CYCLIC_SHELL *ctx;
35		348	const PetscScalar d_px,read;
36		348	PetscScalar d_py,write;
37		348	PetscInt m,n;
38
39		348	PetscFunctionBegin;
40	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 2 times.	348	PetscCall(MatShellGetContext(B,&ctx));
41	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 2 times.	348	PetscCall(MatGetLocalSize(ctx->A,&m,&n));
42	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 2 times.	348	PetscCall(VecCUDAGetArrayRead(x,&d_px));
43	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 2 times.	348	PetscCall(VecCUDAGetArrayWrite(y,&d_py));
44			/* dense matrix columns may misalign either block; retain the same Vec objects on every rank */
45	2/2 ✓ Branch 0 taken 1 times. ✓ Branch 1 taken 2 times.	348	if (((PETSC_UINTPTR_T)(d_px))%16) {
46	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 1 times.	6	PetscCall(VecCUDAGetArrayWrite(ctx->x1,&write));
47	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 1 times.	6	PetscCall(Copy_Cyclic_CUDA(write,d_px,m));
48	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 1 times.	6	PetscCall(VecCUDARestoreArrayWrite(ctx->x1,&write));
49	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 2 times.	342	} else PetscCall(VecCUDAPlaceArray(ctx->x1,d_px));
50	2/2 ✓ Branch 0 taken 1 times. ✓ Branch 1 taken 2 times.	348	if (((PETSC_UINTPTR_T)(d_px+m))%16) {
51	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 1 times.	167	PetscCall(VecCUDAGetArrayWrite(ctx->x2,&write));
52	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 1 times.	167	PetscCall(Copy_Cyclic_CUDA(write,d_px+m,n));
53	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 1 times.	167	PetscCall(VecCUDARestoreArrayWrite(ctx->x2,&write));
54	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 2 times.	181	} else PetscCall(VecCUDAPlaceArray(ctx->x2,d_px+m));
55	3/4 ✓ Branch 0 taken 2 times. ✓ Branch 1 taken 1 times. ✗ Branch 2 not taken. ✓ Branch 3 taken 2 times.	348	if (!(((PETSC_UINTPTR_T)(d_py))%16)) PetscCall(VecCUDAPlaceArray(ctx->y1,d_py));
56	3/4 ✓ Branch 0 taken 2 times. ✓ Branch 1 taken 1 times. ✗ Branch 2 not taken. ✓ Branch 3 taken 2 times.	348	if (!(((PETSC_UINTPTR_T)(d_py+m))%16)) PetscCall(VecCUDAPlaceArray(ctx->y2,d_py+m));
57	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 2 times.	348	PetscCall(MatMult(ctx->A,ctx->x2,ctx->y1));
58	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 2 times.	348	PetscCall(MatMult(ctx->AT,ctx->x1,ctx->y2));
59	2/2 ✓ Branch 0 taken 1 times. ✓ Branch 1 taken 2 times.	348	if (((PETSC_UINTPTR_T)(d_py))%16) {
60	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 1 times.	6	PetscCall(VecCUDAGetArrayRead(ctx->y1,&read));
61	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 1 times.	6	PetscCall(Copy_Cyclic_CUDA(d_py,read,m));
62	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 1 times.	6	PetscCall(VecCUDARestoreArrayRead(ctx->y1,&read));
63	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 2 times.	342	} else PetscCall(VecCUDAResetArray(ctx->y1));
64	2/2 ✓ Branch 0 taken 1 times. ✓ Branch 1 taken 2 times.	348	if (((PETSC_UINTPTR_T)(d_py+m))%16) {
65	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 1 times.	167	PetscCall(VecCUDAGetArrayRead(ctx->y2,&read));
66	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 1 times.	167	PetscCall(Copy_Cyclic_CUDA(d_py+m,read,n));
67	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 1 times.	167	PetscCall(VecCUDARestoreArrayRead(ctx->y2,&read));
68	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 2 times.	181	} else PetscCall(VecCUDAResetArray(ctx->y2));
69	3/4 ✓ Branch 0 taken 2 times. ✓ Branch 1 taken 1 times. ✗ Branch 2 not taken. ✓ Branch 3 taken 2 times.	348	if (!(((PETSC_UINTPTR_T)(d_px))%16)) PetscCall(VecCUDAResetArray(ctx->x1));
70	3/4 ✓ Branch 0 taken 2 times. ✓ Branch 1 taken 1 times. ✗ Branch 2 not taken. ✓ Branch 3 taken 2 times.	348	if (!(((PETSC_UINTPTR_T)(d_px+m))%16)) PetscCall(VecCUDAResetArray(ctx->x2));
71	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 2 times.	348	PetscCall(VecCUDARestoreArrayRead(x,&d_px));
72	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 2 times.	348	PetscCall(VecCUDARestoreArrayWrite(y,&d_py));
73			PetscFunctionReturn(PETSC_SUCCESS);
74			}
75
76		4861	PetscErrorCode MatMult_ECross_CUDA(Mat B,Vec x,Vec y)
77			{
78		4861	SVD_CYCLIC_SHELL *ctx;
79		4861	const PetscScalar d_px,read;
80		4861	PetscScalar d_py,write;
81		4861	PetscInt mn,m,n;
82
83		4861	PetscFunctionBegin;
84	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 2 times.	4861	PetscCall(MatShellGetContext(B,&ctx));
85	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 2 times.	4861	PetscCall(MatGetLocalSize(ctx->A,NULL,&n));
86	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 2 times.	4861	PetscCall(VecGetLocalSize(y,&mn));
87		4861	m = mn-n;
88	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 2 times.	4861	PetscCall(VecCUDAGetArrayRead(x,&d_px));
89	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 2 times.	4861	PetscCall(VecCUDAGetArrayWrite(y,&d_py));
90	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 2 times.	4861	PetscCall(VecCUDAPlaceArray(ctx->x1,d_px));
91	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 2 times.	4861	PetscCall(VecCUDAPlaceArray(ctx->y1,d_py));
92	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 2 times.	4861	PetscCall(VecCopy(ctx->x1,ctx->y1));
93	2/2 ✓ Branch 0 taken 1 times. ✓ Branch 1 taken 1 times.	4861	if (((PETSC_UINTPTR_T)(d_px+m))%16) {
94	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 1 times.	2408	PetscCall(VecCUDAGetArrayWrite(ctx->x2,&write));
95	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 1 times.	2408	PetscCall(Copy_Cyclic_CUDA(write,d_px+m,n));
96	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 1 times.	2408	PetscCall(VecCUDARestoreArrayWrite(ctx->x2,&write));
97	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 1 times.	2453	} else PetscCall(VecCUDAPlaceArray(ctx->x2,d_px+m));
98	3/4 ✓ Branch 0 taken 1 times. ✓ Branch 1 taken 1 times. ✗ Branch 2 not taken. ✓ Branch 3 taken 1 times.	4861	if (!(((PETSC_UINTPTR_T)(d_py+m))%16)) PetscCall(VecCUDAPlaceArray(ctx->y2,d_py+m));
99	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 2 times.	4861	PetscCall(MatMult(ctx->A,ctx->x2,ctx->w));
100	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 2 times.	4861	PetscCall(MatMult(ctx->AT,ctx->w,ctx->y2));
101	2/2 ✓ Branch 0 taken 1 times. ✓ Branch 1 taken 1 times.	4861	if (((PETSC_UINTPTR_T)(d_py+m))%16) {
102	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 1 times.	2408	PetscCall(VecCUDAGetArrayRead(ctx->y2,&read));
103	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 1 times.	2408	PetscCall(Copy_Cyclic_CUDA(d_py+m,read,n));
104	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 1 times.	2408	PetscCall(VecCUDARestoreArrayRead(ctx->y2,&read));
105	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 1 times.	2453	} else PetscCall(VecCUDAResetArray(ctx->y2));
106	3/4 ✓ Branch 0 taken 1 times. ✓ Branch 1 taken 1 times. ✗ Branch 2 not taken. ✓ Branch 3 taken 1 times.	4861	if (!(((PETSC_UINTPTR_T)(d_px+m))%16)) PetscCall(VecCUDAResetArray(ctx->x2));
107	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 2 times.	4861	PetscCall(VecCUDAResetArray(ctx->x1));
108	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 2 times.	4861	PetscCall(VecCUDAResetArray(ctx->y1));
109	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 2 times.	4861	PetscCall(VecCUDARestoreArrayRead(x,&d_px));
110	1/2 ✗ Branch 0 not taken. ✓ Branch 1 taken 2 times.	4861	PetscCall(VecCUDARestoreArrayWrite(y,&d_py));
111			PetscFunctionReturn(PETSC_SUCCESS);
112			}
113

Function (Line)	Call count	Block coverage
MatMult_Cyclic_CUDA (line 32)	called 348 times	70.0%
MatMult_ECross_CUDA (line 76)	called 4861 times	58.0%
_ZL16Copy_Cyclic_CUDAPN6thrust16_V_300403_SM_8607complexIdEEPKS2_i (line 18)	not called	0.0%
_ZL16Copy_Cyclic_CUDAPdPKdi (line 18)	called 5162 times	55.0%