| 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 | CUDA-related code common to several BV impls | ||
| 12 | */ | ||
| 13 | |||
| 14 | #include <slepc/private/bvimpl.h> | ||
| 15 | #include <slepccupmblas.h> | ||
| 16 | |||
| 17 | #define BLOCKSIZE 64 | ||
| 18 | |||
| 19 | /* | ||
| 20 | C := alpha*A*B + beta*C | ||
| 21 | */ | ||
| 22 | 496 | PetscErrorCode BVMult_BLAS_CUDA(BV,PetscInt m_,PetscInt n_,PetscInt k_,PetscScalar alpha,const PetscScalar *d_A,PetscInt lda_,const PetscScalar *d_B,PetscInt ldb_,PetscScalar beta,PetscScalar *d_C,PetscInt ldc_) | |
| 23 | { | ||
| 24 | 496 | PetscCuBLASInt m=0,n=0,k=0,lda=0,ldb=0,ldc=0; | |
| 25 | 496 | cublasHandle_t cublasv2handle; | |
| 26 | |||
| 27 | 496 | PetscFunctionBegin; | |
| 28 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
496 | PetscCall(PetscCUBLASGetHandle(&cublasv2handle)); |
| 29 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
496 | PetscCall(PetscCuBLASIntCast(m_,&m)); |
| 30 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
496 | PetscCall(PetscCuBLASIntCast(n_,&n)); |
| 31 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
496 | PetscCall(PetscCuBLASIntCast(k_,&k)); |
| 32 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
496 | PetscCall(PetscCuBLASIntCast(lda_,&lda)); |
| 33 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
496 | PetscCall(PetscCuBLASIntCast(ldb_,&ldb)); |
| 34 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
496 | PetscCall(PetscCuBLASIntCast(ldc_,&ldc)); |
| 35 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
496 | PetscCall(PetscLogGpuTimeBegin()); |
| 36 |
1/10✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
✗ Branch 4 not taken.
✗ Branch 5 not taken.
✗ Branch 6 not taken.
✗ Branch 7 not taken.
✗ Branch 8 not taken.
✗ Branch 9 not taken.
|
496 | PetscCallCUBLAS(cublasXgemm(cublasv2handle,CUBLAS_OP_N,CUBLAS_OP_N,m,n,k,&alpha,d_A,lda,d_B,ldb,&beta,d_C,ldc)); |
| 37 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
496 | PetscCall(PetscLogGpuTimeEnd()); |
| 38 | 496 | PetscCall(PetscLogGpuFlops(2.0*m*n*k)); | |
| 39 | 496 | PetscFunctionReturn(PETSC_SUCCESS); | |
| 40 | } | ||
| 41 | |||
| 42 | /* | ||
| 43 | y := alpha*A*x + beta*y | ||
| 44 | */ | ||
| 45 | 78499 | PetscErrorCode BVMultVec_BLAS_CUDA(BV,PetscInt n_,PetscInt k_,PetscScalar alpha,const PetscScalar *d_A,PetscInt lda_,const PetscScalar *d_x,PetscScalar beta,PetscScalar *d_y) | |
| 46 | { | ||
| 47 | 78499 | PetscCuBLASInt n=0,k=0,lda=0,one=1; | |
| 48 | 78499 | cublasHandle_t cublasv2handle; | |
| 49 | |||
| 50 | 78499 | PetscFunctionBegin; | |
| 51 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
78499 | PetscCall(PetscCUBLASGetHandle(&cublasv2handle)); |
| 52 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
78499 | PetscCall(PetscCuBLASIntCast(n_,&n)); |
| 53 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
78499 | PetscCall(PetscCuBLASIntCast(k_,&k)); |
| 54 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
78499 | PetscCall(PetscCuBLASIntCast(lda_,&lda)); |
| 55 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
78499 | PetscCall(PetscLogGpuTimeBegin()); |
| 56 |
1/10✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
✗ Branch 4 not taken.
✗ Branch 5 not taken.
✗ Branch 6 not taken.
✗ Branch 7 not taken.
✗ Branch 8 not taken.
✗ Branch 9 not taken.
|
78499 | PetscCallCUBLAS(cublasXgemv(cublasv2handle,CUBLAS_OP_N,n,k,&alpha,d_A,lda,d_x,one,&beta,d_y,one)); |
| 57 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
78499 | PetscCall(PetscLogGpuTimeEnd()); |
| 58 | 78499 | PetscCall(PetscLogGpuFlops(2.0*n*k)); | |
| 59 | 78499 | PetscFunctionReturn(PETSC_SUCCESS); | |
| 60 | } | ||
| 61 | |||
| 62 | /* | ||
| 63 | A(:,s:e-1) := A*B(:,s:e-1) | ||
| 64 | */ | ||
| 65 | 3833 | PetscErrorCode BVMultInPlace_BLAS_CUDA(BV,PetscInt m_,PetscInt k_,PetscInt s,PetscInt e,PetscScalar *d_A,PetscInt lda_,const PetscScalar *d_B,PetscInt ldb_,PetscBool btrans) | |
| 66 | { | ||
| 67 | 3833 | const PetscScalar *d_B1; | |
| 68 | 3833 | PetscScalar *d_work,sone=1.0,szero=0.0; | |
| 69 | 3833 | PetscCuBLASInt m=0,n=0,k=0,l=0,lda=0,ldb=0,bs=BLOCKSIZE; | |
| 70 | 3833 | size_t freemem,totmem; | |
| 71 | 3833 | cublasHandle_t cublasv2handle; | |
| 72 | 3833 | cublasOperation_t bt; | |
| 73 | |||
| 74 | 3833 | PetscFunctionBegin; | |
| 75 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
3833 | PetscCall(PetscCUBLASGetHandle(&cublasv2handle)); |
| 76 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
3833 | PetscCall(PetscCuBLASIntCast(m_,&m)); |
| 77 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
3833 | PetscCall(PetscCuBLASIntCast(e-s,&n)); |
| 78 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
3833 | PetscCall(PetscCuBLASIntCast(k_,&k)); |
| 79 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
3833 | PetscCall(PetscCuBLASIntCast(lda_,&lda)); |
| 80 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
3833 | PetscCall(PetscCuBLASIntCast(ldb_,&ldb)); |
| 81 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
3833 | PetscCall(PetscLogGpuTimeBegin()); |
| 82 |
2/2✓ Branch 0 taken 2 times.
✓ Branch 1 taken 2 times.
|
3833 | if (PetscUnlikely(btrans)) { |
| 83 | 4 | d_B1 = d_B+s; | |
| 84 | 4 | bt = CUBLAS_OP_C; | |
| 85 | } else { | ||
| 86 | 3829 | d_B1 = d_B+s*ldb; | |
| 87 | 3829 | bt = CUBLAS_OP_N; | |
| 88 | } | ||
| 89 | /* try to allocate the whole matrix */ | ||
| 90 |
1/4✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
|
3833 | PetscCallCUDA(cudaMemGetInfo(&freemem,&totmem)); |
| 91 |
1/2✓ Branch 0 taken 2 times.
✗ Branch 1 not taken.
|
3833 | if (freemem>=lda*n*sizeof(PetscScalar)) { |
| 92 |
1/4✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
|
3833 | PetscCallCUDA(cudaMalloc((void**)&d_work,lda*n*sizeof(PetscScalar))); |
| 93 |
1/10✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
✗ Branch 4 not taken.
✗ Branch 5 not taken.
✗ Branch 6 not taken.
✗ Branch 7 not taken.
✗ Branch 8 not taken.
✗ Branch 9 not taken.
|
3833 | PetscCallCUBLAS(cublasXgemm(cublasv2handle,CUBLAS_OP_N,bt,m,n,k,&sone,d_A,lda,d_B1,ldb,&szero,d_work,lda)); |
| 94 |
1/4✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
|
3833 | PetscCallCUDA(cudaMemcpy2D(d_A+s*lda,lda*sizeof(PetscScalar),d_work,lda*sizeof(PetscScalar),m*sizeof(PetscScalar),n,cudaMemcpyDeviceToDevice)); |
| 95 | } else { | ||
| 96 | ✗ | PetscCall(PetscCuBLASIntCast(freemem/(m*sizeof(PetscScalar)),&bs)); | |
| 97 | ✗ | PetscCallCUDA(cudaMalloc((void**)&d_work,bs*n*sizeof(PetscScalar))); | |
| 98 | ✗ | PetscCall(PetscCuBLASIntCast(m % bs,&l)); | |
| 99 | ✗ | if (l) { | |
| 100 | ✗ | PetscCallCUBLAS(cublasXgemm(cublasv2handle,CUBLAS_OP_N,bt,l,n,k,&sone,d_A,lda,d_B1,ldb,&szero,d_work,l)); | |
| 101 | ✗ | PetscCallCUDA(cudaMemcpy2D(d_A+s*lda,lda*sizeof(PetscScalar),d_work,l*sizeof(PetscScalar),l*sizeof(PetscScalar),n,cudaMemcpyDeviceToDevice)); | |
| 102 | } | ||
| 103 | ✗ | for (;l<m;l+=bs) { | |
| 104 | ✗ | PetscCallCUBLAS(cublasXgemm(cublasv2handle,CUBLAS_OP_N,bt,bs,n,k,&sone,d_A+l,lda,d_B1,ldb,&szero,d_work,bs)); | |
| 105 | ✗ | PetscCallCUDA(cudaMemcpy2D(d_A+l+s*lda,lda*sizeof(PetscScalar),d_work,bs*sizeof(PetscScalar),bs*sizeof(PetscScalar),n,cudaMemcpyDeviceToDevice)); | |
| 106 | } | ||
| 107 | } | ||
| 108 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
3833 | PetscCall(PetscLogGpuTimeEnd()); |
| 109 |
1/4✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
|
3833 | PetscCallCUDA(cudaFree(d_work)); |
| 110 | 3833 | PetscCall(PetscLogGpuFlops(2.0*m*n*k)); | |
| 111 | 3833 | PetscFunctionReturn(PETSC_SUCCESS); | |
| 112 | } | ||
| 113 | |||
| 114 | /* | ||
| 115 | B := alpha*A + beta*B | ||
| 116 | */ | ||
| 117 | 274 | PetscErrorCode BVAXPY_BLAS_CUDA(BV,PetscInt n_,PetscInt k_,PetscScalar alpha,const PetscScalar *d_A,PetscInt lda_,PetscScalar beta,PetscScalar *d_B,PetscInt ldb_) | |
| 118 | { | ||
| 119 | 274 | PetscCuBLASInt n=0,k=0,lda=0,ldb=0; | |
| 120 | 274 | cublasHandle_t cublasv2handle; | |
| 121 | |||
| 122 | 274 | PetscFunctionBegin; | |
| 123 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
274 | PetscCall(PetscCUBLASGetHandle(&cublasv2handle)); |
| 124 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
274 | PetscCall(PetscCuBLASIntCast(n_,&n)); |
| 125 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
274 | PetscCall(PetscCuBLASIntCast(k_,&k)); |
| 126 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
274 | PetscCall(PetscCuBLASIntCast(lda_,&lda)); |
| 127 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
274 | PetscCall(PetscCuBLASIntCast(ldb_,&ldb)); |
| 128 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
274 | PetscCall(PetscLogGpuTimeBegin()); |
| 129 |
1/10✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
✗ Branch 4 not taken.
✗ Branch 5 not taken.
✗ Branch 6 not taken.
✗ Branch 7 not taken.
✗ Branch 8 not taken.
✗ Branch 9 not taken.
|
274 | PetscCallCUBLAS(cublasXgeam(cublasv2handle,CUBLAS_OP_N,CUBLAS_OP_N,n,k,&alpha,d_A,lda,&beta,d_B,ldb,d_B,ldb)); |
| 130 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
274 | PetscCall(PetscLogGpuTimeEnd()); |
| 131 |
1/2✓ Branch 0 taken 2 times.
✗ Branch 1 not taken.
|
415 | PetscCall(PetscLogGpuFlops((beta==(PetscScalar)1.0)?2.0*n*k:3.0*n*k)); |
| 132 | 274 | PetscFunctionReturn(PETSC_SUCCESS); | |
| 133 | } | ||
| 134 | |||
| 135 | /* | ||
| 136 | C := A'*B | ||
| 137 | |||
| 138 | C is a CPU array | ||
| 139 | */ | ||
| 140 | 5330 | PetscErrorCode BVDot_BLAS_CUDA(BV bv,PetscInt m_,PetscInt n_,PetscInt k_,const PetscScalar *d_A,PetscInt lda_,const PetscScalar *d_B,PetscInt ldb_,PetscScalar *C,PetscInt ldc_,PetscBool mpi) | |
| 141 | { | ||
| 142 | 5330 | PetscScalar *d_work,sone=1.0,szero=0.0,*CC; | |
| 143 | 5330 | PetscInt j; | |
| 144 | 5330 | PetscCuBLASInt m=0,n=0,k=0,lda=0,ldb=0,ldc=0; | |
| 145 | 5330 | PetscMPIInt len; | |
| 146 | 5330 | cublasHandle_t cublasv2handle; | |
| 147 | |||
| 148 | 5330 | PetscFunctionBegin; | |
| 149 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
5330 | PetscCall(PetscCUBLASGetHandle(&cublasv2handle)); |
| 150 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
5330 | PetscCall(PetscCuBLASIntCast(m_,&m)); |
| 151 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
5330 | PetscCall(PetscCuBLASIntCast(n_,&n)); |
| 152 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
5330 | PetscCall(PetscCuBLASIntCast(k_,&k)); |
| 153 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
5330 | PetscCall(PetscCuBLASIntCast(lda_,&lda)); |
| 154 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
5330 | PetscCall(PetscCuBLASIntCast(ldb_,&ldb)); |
| 155 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
5330 | PetscCall(PetscCuBLASIntCast(ldc_,&ldc)); |
| 156 |
1/4✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
|
5330 | PetscCallCUDA(cudaMalloc((void**)&d_work,m*n*sizeof(PetscScalar))); |
| 157 |
2/2✓ Branch 0 taken 2 times.
✓ Branch 1 taken 2 times.
|
5330 | if (mpi) { |
| 158 |
2/2✓ Branch 0 taken 2 times.
✓ Branch 1 taken 2 times.
|
608 | if (ldc==m) { |
| 159 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
128 | PetscCall(BVAllocateWork_Private(bv,m*n)); |
| 160 |
1/2✓ Branch 0 taken 2 times.
✗ Branch 1 not taken.
|
128 | if (k) { |
| 161 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
128 | PetscCall(PetscLogGpuTimeBegin()); |
| 162 |
1/10✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
✗ Branch 4 not taken.
✗ Branch 5 not taken.
✗ Branch 6 not taken.
✗ Branch 7 not taken.
✗ Branch 8 not taken.
✗ Branch 9 not taken.
|
128 | PetscCallCUBLAS(cublasXgemm(cublasv2handle,CUBLAS_OP_C,CUBLAS_OP_N,m,n,k,&sone,d_A,lda,d_B,ldb,&szero,d_work,ldc)); |
| 163 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
128 | PetscCall(PetscLogGpuTimeEnd()); |
| 164 |
1/4✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
|
128 | PetscCallCUDA(cudaMemcpy(bv->work,d_work,m*n*sizeof(PetscScalar),cudaMemcpyDeviceToHost)); |
| 165 | 128 | PetscCall(PetscLogGpuToCpu(m*n*sizeof(PetscScalar))); | |
| 166 | ✗ | } else PetscCall(PetscArrayzero(bv->work,m*n)); | |
| 167 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
128 | PetscCall(PetscMPIIntCast(m*n,&len)); |
| 168 |
2/6✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✓ Branch 3 taken 2 times.
✗ Branch 4 not taken.
✗ Branch 5 not taken.
|
128 | PetscCallMPI(MPIU_Allreduce(bv->work,C,len,MPIU_SCALAR,MPIU_SUM,PetscObjectComm((PetscObject)bv))); |
| 169 | } else { | ||
| 170 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
480 | PetscCall(BVAllocateWork_Private(bv,2*m*n)); |
| 171 | 480 | CC = bv->work+m*n; | |
| 172 |
1/2✓ Branch 0 taken 2 times.
✗ Branch 1 not taken.
|
480 | if (k) { |
| 173 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
480 | PetscCall(PetscLogGpuTimeBegin()); |
| 174 |
1/10✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
✗ Branch 4 not taken.
✗ Branch 5 not taken.
✗ Branch 6 not taken.
✗ Branch 7 not taken.
✗ Branch 8 not taken.
✗ Branch 9 not taken.
|
480 | PetscCallCUBLAS(cublasXgemm(cublasv2handle,CUBLAS_OP_C,CUBLAS_OP_N,m,n,k,&sone,d_A,lda,d_B,ldb,&szero,d_work,m)); |
| 175 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
480 | PetscCall(PetscLogGpuTimeEnd()); |
| 176 |
1/4✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
|
480 | PetscCallCUDA(cudaMemcpy(bv->work,d_work,m*n*sizeof(PetscScalar),cudaMemcpyDeviceToHost)); |
| 177 | 480 | PetscCall(PetscLogGpuToCpu(m*n*sizeof(PetscScalar))); | |
| 178 | ✗ | } else PetscCall(PetscArrayzero(bv->work,m*n)); | |
| 179 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
480 | PetscCall(PetscMPIIntCast(m*n,&len)); |
| 180 |
2/6✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✓ Branch 3 taken 2 times.
✗ Branch 4 not taken.
✗ Branch 5 not taken.
|
480 | PetscCallMPI(MPIU_Allreduce(bv->work,CC,len,MPIU_SCALAR,MPIU_SUM,PetscObjectComm((PetscObject)bv))); |
| 181 |
3/4✓ Branch 0 taken 2 times.
✗ Branch 1 not taken.
✓ Branch 2 taken 2 times.
✓ Branch 3 taken 2 times.
|
2592 | for (j=0;j<n;j++) PetscCall(PetscArraycpy(C+j*ldc,CC+j*m,m)); |
| 182 | } | ||
| 183 | } else { | ||
| 184 |
1/2✓ Branch 0 taken 2 times.
✗ Branch 1 not taken.
|
4722 | if (k) { |
| 185 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
4722 | PetscCall(BVAllocateWork_Private(bv,m*n)); |
| 186 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
4722 | PetscCall(PetscLogGpuTimeBegin()); |
| 187 |
1/10✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
✗ Branch 4 not taken.
✗ Branch 5 not taken.
✗ Branch 6 not taken.
✗ Branch 7 not taken.
✗ Branch 8 not taken.
✗ Branch 9 not taken.
|
4722 | PetscCallCUBLAS(cublasXgemm(cublasv2handle,CUBLAS_OP_C,CUBLAS_OP_N,m,n,k,&sone,d_A,lda,d_B,ldb,&szero,d_work,m)); |
| 188 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
4722 | PetscCall(PetscLogGpuTimeEnd()); |
| 189 |
1/4✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
|
4722 | PetscCallCUDA(cudaMemcpy(bv->work,d_work,m*n*sizeof(PetscScalar),cudaMemcpyDeviceToHost)); |
| 190 | 4722 | PetscCall(PetscLogGpuToCpu(m*n*sizeof(PetscScalar))); | |
| 191 |
3/4✓ Branch 0 taken 2 times.
✗ Branch 1 not taken.
✓ Branch 2 taken 2 times.
✓ Branch 3 taken 2 times.
|
32016 | for (j=0;j<n;j++) PetscCall(PetscArraycpy(C+j*ldc,bv->work+j*m,m)); |
| 192 | } | ||
| 193 | } | ||
| 194 |
1/4✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
|
5330 | PetscCallCUDA(cudaFree(d_work)); |
| 195 | 5330 | PetscCall(PetscLogGpuFlops(2.0*m*n*k)); | |
| 196 | 5330 | PetscFunctionReturn(PETSC_SUCCESS); | |
| 197 | } | ||
| 198 | |||
| 199 | /* | ||
| 200 | y := A'*x | ||
| 201 | |||
| 202 | y is a CPU array, if NULL bv->buffer is used as a workspace | ||
| 203 | */ | ||
| 204 | 49983 | PetscErrorCode BVDotVec_BLAS_CUDA(BV bv,PetscInt n_,PetscInt k_,const PetscScalar *d_A,PetscInt lda_,const PetscScalar *d_x,PetscScalar *y,PetscBool mpi) | |
| 205 | { | ||
| 206 | 49983 | PetscScalar *d_work,szero=0.0,sone=1.0,*yy; | |
| 207 | 49983 | PetscCuBLASInt n=0,k=0,lda=0,one=1; | |
| 208 | 49983 | PetscMPIInt len; | |
| 209 | 49983 | cublasHandle_t cublasv2handle; | |
| 210 | |||
| 211 | 49983 | PetscFunctionBegin; | |
| 212 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
49983 | PetscCall(PetscCUBLASGetHandle(&cublasv2handle)); |
| 213 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
49983 | PetscCall(PetscCuBLASIntCast(n_,&n)); |
| 214 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
49983 | PetscCall(PetscCuBLASIntCast(k_,&k)); |
| 215 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
49983 | PetscCall(PetscCuBLASIntCast(lda_,&lda)); |
| 216 |
3/4✓ Branch 0 taken 2 times.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✓ Branch 3 taken 2 times.
|
49983 | if (!y) PetscCall(VecGetArrayWriteAndMemType(bv->buffer,&d_work,NULL)); |
| 217 |
1/4✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
|
12620 | else PetscCallCUDA(cudaMalloc((void**)&d_work,k*sizeof(PetscScalar))); |
| 218 |
2/2✓ Branch 0 taken 2 times.
✓ Branch 1 taken 2 times.
|
49983 | if (mpi) { |
| 219 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
1952 | PetscCall(BVAllocateWork_Private(bv,k)); |
| 220 |
1/2✓ Branch 0 taken 2 times.
✗ Branch 1 not taken.
|
1952 | if (n) { |
| 221 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
1952 | PetscCall(PetscLogGpuTimeBegin()); |
| 222 |
1/10✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
✗ Branch 4 not taken.
✗ Branch 5 not taken.
✗ Branch 6 not taken.
✗ Branch 7 not taken.
✗ Branch 8 not taken.
✗ Branch 9 not taken.
|
1952 | PetscCallCUBLAS(cublasXgemv(cublasv2handle,CUBLAS_OP_C,n,k,&sone,d_A,lda,d_x,one,&szero,d_work,one)); |
| 223 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
1952 | PetscCall(PetscLogGpuTimeEnd()); |
| 224 |
1/4✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
|
1952 | PetscCallCUDA(cudaMemcpy(bv->work,d_work,k*sizeof(PetscScalar),cudaMemcpyDeviceToHost)); |
| 225 | 1952 | PetscCall(PetscLogGpuToCpu(k*sizeof(PetscScalar))); | |
| 226 | ✗ | } else PetscCall(PetscArrayzero(bv->work,k)); | |
| 227 | /* reduction */ | ||
| 228 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
1952 | PetscCall(PetscMPIIntCast(k,&len)); |
| 229 |
2/2✓ Branch 0 taken 2 times.
✓ Branch 1 taken 2 times.
|
1952 | if (!y) { |
| 230 |
1/2✓ Branch 0 taken 2 times.
✗ Branch 1 not taken.
|
528 | if (use_gpu_aware_mpi) { /* case 1: reduce on GPU using a temporary buffer */ |
| 231 |
1/4✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
|
528 | PetscCallCUDA(cudaMalloc((void**)&yy,k*sizeof(PetscScalar))); |
| 232 |
2/6✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✓ Branch 3 taken 2 times.
✗ Branch 4 not taken.
✗ Branch 5 not taken.
|
528 | PetscCallMPI(MPIU_Allreduce(d_work,yy,len,MPIU_SCALAR,MPIU_SUM,PetscObjectComm((PetscObject)bv))); |
| 233 |
1/4✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
|
528 | PetscCallCUDA(cudaMemcpy(d_work,yy,k*sizeof(PetscScalar),cudaMemcpyDeviceToDevice)); |
| 234 |
1/4✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
|
528 | PetscCallCUDA(cudaFree(yy)); |
| 235 | } else { /* case 2: reduce on CPU, copy result back to GPU */ | ||
| 236 | ✗ | PetscCall(BVAllocateWork_Private(bv,2*k)); | |
| 237 | ✗ | yy = bv->work+k; | |
| 238 | ✗ | PetscCallCUDA(cudaMemcpy(bv->work,d_work,k*sizeof(PetscScalar),cudaMemcpyDeviceToHost)); | |
| 239 | ✗ | PetscCall(PetscLogGpuToCpu(k*sizeof(PetscScalar))); | |
| 240 | ✗ | PetscCallMPI(MPIU_Allreduce(bv->work,yy,len,MPIU_SCALAR,MPIU_SUM,PetscObjectComm((PetscObject)bv))); | |
| 241 | ✗ | PetscCallCUDA(cudaMemcpy(d_work,yy,k*sizeof(PetscScalar),cudaMemcpyHostToDevice)); | |
| 242 | ✗ | PetscCall(PetscLogCpuToGpu(k*sizeof(PetscScalar))); | |
| 243 | } | ||
| 244 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
528 | PetscCall(VecRestoreArrayWriteAndMemType(bv->buffer,&d_work)); |
| 245 | } else { /* case 3: user-provided array y, reduce on CPU */ | ||
| 246 |
1/4✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
|
1424 | PetscCallCUDA(cudaFree(d_work)); |
| 247 |
2/6✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✓ Branch 3 taken 2 times.
✗ Branch 4 not taken.
✗ Branch 5 not taken.
|
1424 | PetscCallMPI(MPIU_Allreduce(bv->work,y,len,MPIU_SCALAR,MPIU_SUM,PetscObjectComm((PetscObject)bv))); |
| 248 | } | ||
| 249 | } else { | ||
| 250 |
1/2✓ Branch 0 taken 2 times.
✗ Branch 1 not taken.
|
48031 | if (n) { |
| 251 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
48031 | PetscCall(PetscLogGpuTimeBegin()); |
| 252 |
1/10✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
✗ Branch 4 not taken.
✗ Branch 5 not taken.
✗ Branch 6 not taken.
✗ Branch 7 not taken.
✗ Branch 8 not taken.
✗ Branch 9 not taken.
|
48031 | PetscCallCUBLAS(cublasXgemv(cublasv2handle,CUBLAS_OP_C,n,k,&sone,d_A,lda,d_x,one,&szero,d_work,one)); |
| 253 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
48031 | PetscCall(PetscLogGpuTimeEnd()); |
| 254 | } | ||
| 255 |
3/4✓ Branch 0 taken 2 times.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✓ Branch 3 taken 2 times.
|
48031 | if (!y) PetscCall(VecRestoreArrayWriteAndMemType(bv->buffer,&d_work)); |
| 256 | else { | ||
| 257 |
1/4✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
|
11196 | PetscCallCUDA(cudaMemcpy(y,d_work,k*sizeof(PetscScalar),cudaMemcpyDeviceToHost)); |
| 258 | 11196 | PetscCall(PetscLogGpuToCpu(k*sizeof(PetscScalar))); | |
| 259 |
1/4✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
|
11196 | PetscCallCUDA(cudaFree(d_work)); |
| 260 | } | ||
| 261 | } | ||
| 262 | 49983 | PetscCall(PetscLogGpuFlops(2.0*n*k)); | |
| 263 | 49983 | PetscFunctionReturn(PETSC_SUCCESS); | |
| 264 | } | ||
| 265 | |||
| 266 | /* | ||
| 267 | Scale n scalars | ||
| 268 | */ | ||
| 269 | 27949 | PetscErrorCode BVScale_BLAS_CUDA(BV,PetscInt n_,PetscScalar *d_A,PetscScalar alpha) | |
| 270 | { | ||
| 271 | 27949 | PetscCuBLASInt n=0,one=1; | |
| 272 | 27949 | cublasHandle_t cublasv2handle; | |
| 273 | |||
| 274 | 27949 | PetscFunctionBegin; | |
| 275 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
27949 | PetscCall(PetscCuBLASIntCast(n_,&n)); |
| 276 |
5/8✓ Branch 0 taken 2 times.
✓ Branch 1 taken 2 times.
✓ Branch 2 taken 1 times.
✓ Branch 3 taken 2 times.
✗ Branch 4 not taken.
✓ Branch 5 taken 1 times.
✗ Branch 6 not taken.
✗ Branch 7 not taken.
|
41666 | if (PetscUnlikely(alpha == (PetscScalar)0.0)) PetscCallCUDA(cudaMemset(d_A,0,n*sizeof(PetscScalar))); |
| 277 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
41456 | else if (alpha != (PetscScalar)1.0) { |
| 278 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
27739 | PetscCall(PetscCUBLASGetHandle(&cublasv2handle)); |
| 279 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
27739 | PetscCall(PetscLogGpuTimeBegin()); |
| 280 |
1/10✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
✗ Branch 4 not taken.
✗ Branch 5 not taken.
✗ Branch 6 not taken.
✗ Branch 7 not taken.
✗ Branch 8 not taken.
✗ Branch 9 not taken.
|
27739 | PetscCallCUBLAS(cublasXscal(cublasv2handle,n,&alpha,d_A,one)); |
| 281 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
27739 | PetscCall(PetscLogGpuTimeEnd()); |
| 282 | 27739 | PetscCall(PetscLogGpuFlops(1.0*n)); | |
| 283 | } | ||
| 284 | PetscFunctionReturn(PETSC_SUCCESS); | ||
| 285 | } | ||
| 286 | |||
| 287 | /* | ||
| 288 | Compute 2-norm of vector consisting of n scalars | ||
| 289 | */ | ||
| 290 | 2612 | PetscErrorCode BVNorm_BLAS_CUDA(BV,PetscInt n_,const PetscScalar *d_A,PetscReal *nrm) | |
| 291 | { | ||
| 292 | 2612 | PetscCuBLASInt n=0,one=1; | |
| 293 | 2612 | cublasHandle_t cublasv2handle; | |
| 294 | |||
| 295 | 2612 | PetscFunctionBegin; | |
| 296 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
2612 | PetscCall(PetscCuBLASIntCast(n_,&n)); |
| 297 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
2612 | PetscCall(PetscCUBLASGetHandle(&cublasv2handle)); |
| 298 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
2612 | PetscCall(PetscLogGpuTimeBegin()); |
| 299 |
1/10✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
✗ Branch 4 not taken.
✗ Branch 5 not taken.
✗ Branch 6 not taken.
✗ Branch 7 not taken.
✗ Branch 8 not taken.
✗ Branch 9 not taken.
|
2612 | PetscCallCUBLAS(cublasXnrm2(cublasv2handle,n,d_A,one,nrm)); |
| 300 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
2612 | PetscCall(PetscLogGpuTimeEnd()); |
| 301 | 2612 | PetscCall(PetscLogGpuFlops(2.0*n)); | |
| 302 | 2612 | PetscFunctionReturn(PETSC_SUCCESS); | |
| 303 | } | ||
| 304 | |||
| 305 | /* | ||
| 306 | Normalize the columns of A | ||
| 307 | */ | ||
| 308 | 512 | PetscErrorCode BVNormalize_BLAS_CUDA(BV,PetscInt m_,PetscInt n_,PetscScalar *d_A,PetscInt lda_,PetscScalar *eigi) | |
| 309 | { | ||
| 310 | 512 | PetscInt j,k; | |
| 311 | 512 | PetscReal nrm,nrm1; | |
| 312 | 512 | PetscScalar alpha; | |
| 313 | 512 | PetscCuBLASInt m=0,one=1; | |
| 314 | 512 | cublasHandle_t cublasv2handle; | |
| 315 | |||
| 316 | 512 | PetscFunctionBegin; | |
| 317 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
512 | PetscCall(PetscCuBLASIntCast(m_,&m)); |
| 318 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
512 | PetscCall(PetscCUBLASGetHandle(&cublasv2handle)); |
| 319 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
512 | PetscCall(PetscLogGpuTimeBegin()); |
| 320 |
2/2✓ Branch 0 taken 2 times.
✓ Branch 1 taken 2 times.
|
9618 | for (j=0;j<n_;j++) { |
| 321 | 9106 | k = 1; | |
| 322 |
4/4✓ Branch 0 taken 1 times.
✓ Branch 1 taken 1 times.
✓ Branch 2 taken 1 times.
✓ Branch 3 taken 1 times.
|
9106 | if (!PetscDefined(USE_COMPLEX) && eigi && eigi[j] != 0.0) k = 2; |
| 323 |
1/10✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
✗ Branch 4 not taken.
✗ Branch 5 not taken.
✗ Branch 6 not taken.
✗ Branch 7 not taken.
✗ Branch 8 not taken.
✗ Branch 9 not taken.
|
9106 | PetscCallCUBLAS(cublasXnrm2(cublasv2handle,m,d_A+j*lda_,one,&nrm)); |
| 324 |
2/2✓ Branch 0 taken 1 times.
✓ Branch 1 taken 1 times.
|
9106 | if (k==2) { |
| 325 |
1/10✗ Branch 0 not taken.
✓ Branch 1 taken 1 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
✗ Branch 4 not taken.
✗ Branch 5 not taken.
✗ Branch 6 not taken.
✗ Branch 7 not taken.
✗ Branch 8 not taken.
✗ Branch 9 not taken.
|
4 | PetscCallCUBLAS(cublasXnrm2(cublasv2handle,m,d_A+(j+1)*lda_,one,&nrm1)); |
| 326 | 4 | nrm = SlepcAbs(nrm,nrm1); | |
| 327 | } | ||
| 328 | 9106 | alpha = 1.0/nrm; | |
| 329 |
1/10✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
✗ Branch 4 not taken.
✗ Branch 5 not taken.
✗ Branch 6 not taken.
✗ Branch 7 not taken.
✗ Branch 8 not taken.
✗ Branch 9 not taken.
|
9106 | PetscCallCUBLAS(cublasXscal(cublasv2handle,m,&alpha,d_A+j*lda_,one)); |
| 330 |
2/2✓ Branch 0 taken 1 times.
✓ Branch 1 taken 1 times.
|
9106 | if (k==2) { |
| 331 |
1/10✗ Branch 0 not taken.
✓ Branch 1 taken 1 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
✗ Branch 4 not taken.
✗ Branch 5 not taken.
✗ Branch 6 not taken.
✗ Branch 7 not taken.
✗ Branch 8 not taken.
✗ Branch 9 not taken.
|
4 | PetscCallCUBLAS(cublasXscal(cublasv2handle,m,&alpha,d_A+(j+1)*lda_,one)); |
| 332 | j++; | ||
| 333 | } | ||
| 334 | } | ||
| 335 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
512 | PetscCall(PetscLogGpuTimeEnd()); |
| 336 | 512 | PetscCall(PetscLogGpuFlops(3.0*m*n_)); | |
| 337 | 512 | PetscFunctionReturn(PETSC_SUCCESS); | |
| 338 | } | ||
| 339 | |||
| 340 | /* | ||
| 341 | BV_CleanCoefficients_CUDA - Sets to zero all entries of column j of the bv buffer | ||
| 342 | */ | ||
| 343 | 29377 | PetscErrorCode BV_CleanCoefficients_CUDA(BV bv,PetscInt j,PetscScalar *h) | |
| 344 | { | ||
| 345 | 29377 | PetscScalar *d_hh,*d_a; | |
| 346 | 29377 | PetscInt i; | |
| 347 | |||
| 348 | 29377 | PetscFunctionBegin; | |
| 349 |
2/2✓ Branch 0 taken 2 times.
✓ Branch 1 taken 2 times.
|
29377 | if (!h) { |
| 350 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
25006 | PetscCall(VecGetArrayAndMemType(bv->buffer,&d_a,NULL)); |
| 351 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
25006 | PetscCall(PetscLogGpuTimeBegin()); |
| 352 | 25006 | d_hh = d_a + j*(bv->nc+bv->m); | |
| 353 |
1/4✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
|
25006 | PetscCallCUDA(cudaMemset(d_hh,0,(bv->nc+j)*sizeof(PetscScalar))); |
| 354 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
25006 | PetscCall(PetscLogGpuTimeEnd()); |
| 355 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
25006 | PetscCall(VecRestoreArrayAndMemType(bv->buffer,&d_a)); |
| 356 | } else { /* cpu memory */ | ||
| 357 |
2/2✓ Branch 0 taken 2 times.
✓ Branch 1 taken 2 times.
|
22505 | for (i=0;i<bv->nc+j;i++) h[i] = 0.0; |
| 358 | } | ||
| 359 | PetscFunctionReturn(PETSC_SUCCESS); | ||
| 360 | } | ||
| 361 | |||
| 362 | /* | ||
| 363 | BV_AddCoefficients_CUDA - Add the contents of the scratch (0-th column) of the bv buffer | ||
| 364 | into column j of the bv buffer | ||
| 365 | */ | ||
| 366 | 42577 | PetscErrorCode BV_AddCoefficients_CUDA(BV bv,PetscInt j,PetscScalar *h,PetscScalar *c) | |
| 367 | { | ||
| 368 |
2/2✓ Branch 0 taken 1 times.
✓ Branch 1 taken 1 times.
|
42577 | PetscScalar *d_h,*d_c,sone=1.0; |
| 369 | 42577 | PetscInt i; | |
| 370 | 42577 | PetscCuBLASInt idx=0,one=1; | |
| 371 | 42577 | cublasHandle_t cublasv2handle; | |
| 372 | |||
| 373 | 42577 | PetscFunctionBegin; | |
| 374 |
2/2✓ Branch 0 taken 2 times.
✓ Branch 1 taken 2 times.
|
42577 | if (!h) { |
| 375 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
37815 | PetscCall(PetscCUBLASGetHandle(&cublasv2handle)); |
| 376 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
37815 | PetscCall(VecGetArrayAndMemType(bv->buffer,&d_c,NULL)); |
| 377 | 37815 | d_h = d_c + j*(bv->nc+bv->m); | |
| 378 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
37815 | PetscCall(PetscCuBLASIntCast(bv->nc+j,&idx)); |
| 379 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
37815 | PetscCall(PetscLogGpuTimeBegin()); |
| 380 |
1/10✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
✗ Branch 4 not taken.
✗ Branch 5 not taken.
✗ Branch 6 not taken.
✗ Branch 7 not taken.
✗ Branch 8 not taken.
✗ Branch 9 not taken.
|
37815 | PetscCallCUBLAS(cublasXaxpy(cublasv2handle,idx,&sone,d_c,one,d_h,one)); |
| 381 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
37815 | PetscCall(PetscLogGpuTimeEnd()); |
| 382 | 37815 | PetscCall(PetscLogGpuFlops(1.0*(bv->nc+j))); | |
| 383 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
37815 | PetscCall(VecRestoreArrayAndMemType(bv->buffer,&d_c)); |
| 384 | } else { /* cpu memory */ | ||
| 385 |
2/2✓ Branch 0 taken 2 times.
✓ Branch 1 taken 2 times.
|
25122 | for (i=0;i<bv->nc+j;i++) h[i] += c[i]; |
| 386 | 4762 | PetscCall(PetscLogFlops(1.0*(bv->nc+j))); | |
| 387 | } | ||
| 388 | PetscFunctionReturn(PETSC_SUCCESS); | ||
| 389 | } | ||
| 390 | |||
| 391 | /* | ||
| 392 | BV_SetValue_CUDA - Sets value in row j (counted after the constraints) of column k | ||
| 393 | of the coefficients array | ||
| 394 | */ | ||
| 395 | 27870 | PetscErrorCode BV_SetValue_CUDA(BV bv,PetscInt j,PetscInt k,PetscScalar *h,PetscScalar value) | |
| 396 | { | ||
| 397 | 27870 | PetscScalar *d_h,*a; | |
| 398 | |||
| 399 | 27870 | PetscFunctionBegin; | |
| 400 |
2/2✓ Branch 0 taken 2 times.
✓ Branch 1 taken 2 times.
|
27870 | if (!h) { |
| 401 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
27470 | PetscCall(VecGetArrayAndMemType(bv->buffer,&a,NULL)); |
| 402 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
27470 | PetscCall(PetscLogGpuTimeBegin()); |
| 403 | 27470 | d_h = a + k*(bv->nc+bv->m) + bv->nc+j; | |
| 404 |
1/4✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
|
27470 | PetscCallCUDA(cudaMemcpy(d_h,&value,sizeof(PetscScalar),cudaMemcpyHostToDevice)); |
| 405 | 27470 | PetscCall(PetscLogCpuToGpu(sizeof(PetscScalar))); | |
| 406 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
27470 | PetscCall(PetscLogGpuTimeEnd()); |
| 407 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
27470 | PetscCall(VecRestoreArrayAndMemType(bv->buffer,&a)); |
| 408 | } else { /* cpu memory */ | ||
| 409 | 400 | h[bv->nc+j] = value; | |
| 410 | } | ||
| 411 | PetscFunctionReturn(PETSC_SUCCESS); | ||
| 412 | } | ||
| 413 | |||
| 414 | /* | ||
| 415 | BV_SquareSum_CUDA - Returns the value h'*h, where h represents the contents of the | ||
| 416 | coefficients array (up to position j) | ||
| 417 | */ | ||
| 418 | 41677 | PetscErrorCode BV_SquareSum_CUDA(BV bv,PetscInt j,PetscScalar *h,PetscReal *sum) | |
| 419 | { | ||
| 420 | 41677 | const PetscScalar *d_h; | |
| 421 | 41677 | PetscScalar dot; | |
| 422 | 41677 | PetscInt i; | |
| 423 | 41677 | PetscCuBLASInt idx=0,one=1; | |
| 424 | 41677 | cublasHandle_t cublasv2handle; | |
| 425 | |||
| 426 | 41677 | PetscFunctionBegin; | |
| 427 |
2/2✓ Branch 0 taken 2 times.
✓ Branch 1 taken 2 times.
|
41677 | if (!h) { |
| 428 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
37147 | PetscCall(PetscCUBLASGetHandle(&cublasv2handle)); |
| 429 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
37147 | PetscCall(VecGetArrayReadAndMemType(bv->buffer,&d_h,NULL)); |
| 430 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
37147 | PetscCall(PetscCuBLASIntCast(bv->nc+j,&idx)); |
| 431 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
37147 | PetscCall(PetscLogGpuTimeBegin()); |
| 432 |
1/10✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
✗ Branch 4 not taken.
✗ Branch 5 not taken.
✗ Branch 6 not taken.
✗ Branch 7 not taken.
✗ Branch 8 not taken.
✗ Branch 9 not taken.
|
37147 | PetscCallCUBLAS(cublasXdot(cublasv2handle,idx,d_h,one,d_h,one,&dot)); |
| 433 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
37147 | PetscCall(PetscLogGpuTimeEnd()); |
| 434 | 37147 | PetscCall(PetscLogGpuFlops(2.0*(bv->nc+j))); | |
| 435 | 37147 | *sum = PetscRealPart(dot); | |
| 436 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
37147 | PetscCall(VecRestoreArrayReadAndMemType(bv->buffer,&d_h)); |
| 437 | } else { /* cpu memory */ | ||
| 438 | 4530 | *sum = 0.0; | |
| 439 |
2/2✓ Branch 0 taken 2 times.
✓ Branch 1 taken 2 times.
|
24154 | for (i=0;i<bv->nc+j;i++) *sum += PetscRealPart(h[i]*PetscConj(h[i])); |
| 440 | 4530 | PetscCall(PetscLogFlops(2.0*(bv->nc+j))); | |
| 441 | } | ||
| 442 | PetscFunctionReturn(PETSC_SUCCESS); | ||
| 443 | } | ||
| 444 | |||
| 445 | /* pointwise multiplication */ | ||
| 446 | 184 | static __global__ void PointwiseMult_kernel(PetscInt xcount,PetscScalar *a,const PetscScalar *b,PetscInt n) | |
| 447 | ✗ | { | |
| 448 | PetscInt x; | ||
| 449 | |||
| 450 | x = xcount*gridDim.x*blockDim.x+blockIdx.x*blockDim.x+threadIdx.x; | ||
| 451 | if (x<n) a[x] *= PetscRealPart(b[x]); | ||
| 452 | 184 | } | |
| 453 | |||
| 454 | /* pointwise division */ | ||
| 455 | 184 | static __global__ void PointwiseDiv_kernel(PetscInt xcount,PetscScalar *a,const PetscScalar *b,PetscInt n) | |
| 456 | ✗ | { | |
| 457 | PetscInt x; | ||
| 458 | |||
| 459 | x = xcount*gridDim.x*blockDim.x+blockIdx.x*blockDim.x+threadIdx.x; | ||
| 460 | if (x<n) a[x] /= PetscRealPart(b[x]); | ||
| 461 | 184 | } | |
| 462 | |||
| 463 | /* | ||
| 464 | BV_ApplySignature_CUDA - Computes the pointwise product h*omega, where h represents | ||
| 465 | the contents of the coefficients array (up to position j) and omega is the signature; | ||
| 466 | if inverse=TRUE then the operation is h/omega | ||
| 467 | */ | ||
| 468 | 392 | PetscErrorCode BV_ApplySignature_CUDA(BV bv,PetscInt j,PetscScalar *h,PetscBool inverse) | |
| 469 | { | ||
| 470 | 392 | PetscScalar *d_h; | |
| 471 | 392 | const PetscScalar *d_omega,*omega; | |
| 472 | 392 | PetscInt i,xcount; | |
| 473 | 392 | dim3 blocks3d, threads3d; | |
| 474 | |||
| 475 | 392 | PetscFunctionBegin; | |
| 476 |
2/2✓ Branch 0 taken 2 times.
✓ Branch 1 taken 2 times.
|
392 | if (!(bv->nc+j)) PetscFunctionReturn(PETSC_SUCCESS); |
| 477 |
1/2✓ Branch 0 taken 2 times.
✗ Branch 1 not taken.
|
368 | if (!h) { |
| 478 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
368 | PetscCall(VecGetArrayAndMemType(bv->buffer,&d_h,NULL)); |
| 479 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
368 | PetscCall(VecGetArrayReadAndMemType(bv->omega,&d_omega,NULL)); |
| 480 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
368 | PetscCall(SlepcKernelSetGrid1D(bv->nc+j,&blocks3d,&threads3d,&xcount)); |
| 481 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
368 | PetscCall(PetscLogGpuTimeBegin()); |
| 482 |
2/2✓ Branch 0 taken 2 times.
✓ Branch 1 taken 2 times.
|
368 | if (inverse) { |
| 483 |
3/4✓ Branch 0 taken 2 times.
✗ Branch 1 not taken.
✓ Branch 2 taken 2 times.
✓ Branch 3 taken 2 times.
|
368 | for (i=0;i<xcount;i++) PointwiseDiv_kernel<<<blocks3d,threads3d>>>(i,d_h,d_omega,bv->nc+j); |
| 484 | } else { | ||
| 485 |
3/4✓ Branch 0 taken 2 times.
✗ Branch 1 not taken.
✓ Branch 2 taken 2 times.
✓ Branch 3 taken 2 times.
|
368 | for (i=0;i<xcount;i++) PointwiseMult_kernel<<<blocks3d,threads3d>>>(i,d_h,d_omega,bv->nc+j); |
| 486 | } | ||
| 487 |
1/4✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
|
368 | PetscCallCUDA(cudaGetLastError()); |
| 488 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
368 | PetscCall(PetscLogGpuTimeEnd()); |
| 489 | 368 | PetscCall(PetscLogGpuFlops(1.0*(bv->nc+j))); | |
| 490 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
368 | PetscCall(VecRestoreArrayReadAndMemType(bv->omega,&d_omega)); |
| 491 |
1/2✓ Branch 0 taken 2 times.
✗ Branch 1 not taken.
|
368 | PetscCall(VecRestoreArrayAndMemType(bv->buffer,&d_h)); |
| 492 | } else { | ||
| 493 | ✗ | PetscCall(VecGetArrayRead(bv->omega,&omega)); | |
| 494 | ✗ | if (inverse) for (i=0;i<bv->nc+j;i++) h[i] /= PetscRealPart(omega[i]); | |
| 495 | ✗ | else for (i=0;i<bv->nc+j;i++) h[i] *= PetscRealPart(omega[i]); | |
| 496 | ✗ | PetscCall(VecRestoreArrayRead(bv->omega,&omega)); | |
| 497 | ✗ | PetscCall(PetscLogFlops(1.0*(bv->nc+j))); | |
| 498 | } | ||
| 499 | PetscFunctionReturn(PETSC_SUCCESS); | ||
| 500 | } | ||
| 501 | |||
| 502 | /* | ||
| 503 | BV_SquareRoot_CUDA - Returns the square root of position j (counted after the constraints) | ||
| 504 | of the coefficients array | ||
| 505 | */ | ||
| 506 | 37245 | PetscErrorCode BV_SquareRoot_CUDA(BV bv,PetscInt j,PetscScalar *h,PetscReal *beta) | |
| 507 | { | ||
| 508 | 37245 | const PetscScalar *d_h; | |
| 509 | 37245 | PetscScalar hh; | |
| 510 | |||
| 511 | 37245 | PetscFunctionBegin; | |
| 512 |
1/2✓ Branch 0 taken 2 times.
✗ Branch 1 not taken.
|
37245 | if (!h) { |
| 513 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
37245 | PetscCall(VecGetArrayReadAndMemType(bv->buffer,&d_h,NULL)); |
| 514 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
37245 | PetscCall(PetscLogGpuTimeBegin()); |
| 515 |
1/4✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
|
37245 | PetscCallCUDA(cudaMemcpy(&hh,d_h+bv->nc+j,sizeof(PetscScalar),cudaMemcpyDeviceToHost)); |
| 516 | 37245 | PetscCall(PetscLogGpuToCpu(sizeof(PetscScalar))); | |
| 517 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
37245 | PetscCall(PetscLogGpuTimeEnd()); |
| 518 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
37245 | PetscCall(BV_SafeSqrt(bv,hh,beta)); |
| 519 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
37245 | PetscCall(VecRestoreArrayReadAndMemType(bv->buffer,&d_h)); |
| 520 | ✗ | } else PetscCall(BV_SafeSqrt(bv,h[bv->nc+j],beta)); | |
| 521 | PetscFunctionReturn(PETSC_SUCCESS); | ||
| 522 | } | ||
| 523 | |||
| 524 | /* | ||
| 525 | BV_StoreCoefficients_CUDA - Copy the contents of the coefficients array to an array dest | ||
| 526 | provided by the caller (only values from l to j are copied) | ||
| 527 | */ | ||
| 528 | 1353 | PetscErrorCode BV_StoreCoefficients_CUDA(BV bv,PetscInt j,PetscScalar *h,PetscScalar *dest) | |
| 529 | { | ||
| 530 | 1353 | const PetscScalar *d_h,*d_a; | |
| 531 | 1353 | PetscInt i; | |
| 532 | |||
| 533 | 1353 | PetscFunctionBegin; | |
| 534 |
1/2✓ Branch 0 taken 2 times.
✗ Branch 1 not taken.
|
1353 | if (!h) { |
| 535 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
1353 | PetscCall(VecGetArrayReadAndMemType(bv->buffer,&d_a,NULL)); |
| 536 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
1353 | PetscCall(PetscLogGpuTimeBegin()); |
| 537 | 1353 | d_h = d_a + j*(bv->nc+bv->m)+bv->nc; | |
| 538 |
1/4✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
|
1353 | PetscCallCUDA(cudaMemcpy(dest-bv->l,d_h,(j-bv->l)*sizeof(PetscScalar),cudaMemcpyDeviceToHost)); |
| 539 | 1353 | PetscCall(PetscLogGpuToCpu((j-bv->l)*sizeof(PetscScalar))); | |
| 540 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
1353 | PetscCall(PetscLogGpuTimeEnd()); |
| 541 |
1/2✗ Branch 0 not taken.
✓ Branch 1 taken 2 times.
|
1353 | PetscCall(VecRestoreArrayReadAndMemType(bv->buffer,&d_a)); |
| 542 | } else { | ||
| 543 | ✗ | for (i=bv->l;i<j;i++) dest[i-bv->l] = h[bv->nc+i]; | |
| 544 | } | ||
| 545 | PetscFunctionReturn(PETSC_SUCCESS); | ||
| 546 | } | ||
| 547 |