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
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,n*sizeof(*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