GCC Code Coverage Report

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			BV implemented as a single Vec (HIP version)
12			*/
13
14			#include <slepc/private/bvimpl.h>
15			#include <slepccupmblas.h>
16			#include "../src/sys/classes/bv/impls/svec/svec.h"
17
18		✗	PetscErrorCode BVMult_Svec_HIP(BV Y,PetscScalar alpha,PetscScalar beta,BV X,Mat Q)
19			{
20		✗	BV_SVEC *y = (BV_SVEC*)Y->data,*x = (BV_SVEC*)X->data;
21		✗	const PetscScalar *d_px,*d_A,*d_B,*d_q;
22		✗	PetscScalar *d_py,*d_C;
23		✗	PetscInt ldq;
24
25			PetscFunctionBegin;
26		✗	if (!Y->n) PetscFunctionReturn(PETSC_SUCCESS);
27		✗	PetscCall(VecHIPGetArrayRead(x->v,&d_px));
28		✗	if (beta==(PetscScalar)0.0) PetscCall(VecHIPGetArrayWrite(y->v,&d_py));
29		✗	else PetscCall(VecHIPGetArray(y->v,&d_py));
30		✗	d_A = d_px+(X->nc+X->l)*X->ld;
31		✗	d_C = d_py+(Y->nc+Y->l)*Y->ld;
32		✗	if (Q) {
33		✗	PetscCall(MatDenseGetLDA(Q,&ldq));
34		✗	PetscCall(BV_MatDenseHIPGetArrayRead(Y,Q,&d_q));
35		✗	d_B = d_q+Y->l*ldq+X->l;
36		✗	PetscCall(BVMult_BLAS_HIP(Y,Y->n,Y->k-Y->l,X->k-X->l,alpha,d_A,X->ld,d_B,ldq,beta,d_C,Y->ld));
37		✗	PetscCall(BV_MatDenseHIPRestoreArrayRead(Y,Q,&d_q));
38		✗	} else PetscCall(BVAXPY_BLAS_HIP(Y,Y->n,Y->k-Y->l,alpha,d_A,X->ld,beta,d_C,Y->ld));
39		✗	PetscCall(VecHIPRestoreArrayRead(x->v,&d_px));
40		✗	if (beta==(PetscScalar)0.0) PetscCall(VecHIPRestoreArrayWrite(y->v,&d_py));
41		✗	else PetscCall(VecHIPRestoreArray(y->v,&d_py));
42		✗	PetscFunctionReturn(PETSC_SUCCESS);
43		✗	}
44
45		✗	PetscErrorCode BVMultVec_Svec_HIP(BV X,PetscScalar alpha,PetscScalar beta,Vec y,PetscScalar *q)
46			{
47		✗	BV_SVEC *x = (BV_SVEC*)X->data;
48		✗	PetscScalar *d_py,*d_q;
49		✗	const PetscScalar *d_px;
50
51			PetscFunctionBegin;
52		✗	PetscCall(VecHIPGetArrayRead(x->v,&d_px));
53		✗	if (beta==(PetscScalar)0.0) PetscCall(VecHIPGetArrayWrite(y,&d_py));
54		✗	else PetscCall(VecHIPGetArray(y,&d_py));
55		✗	if (!q) PetscCall(VecHIPGetArray(X->buffer,&d_q));
56			else {
57		✗	PetscInt k=X->k-X->l;
58		✗	PetscCallHIP(hipMalloc((void**)&d_q,k*sizeof(PetscScalar)));
59		✗	PetscCallHIP(hipMemcpy(d_q,q,k*sizeof(PetscScalar),hipMemcpyHostToDevice));
60		✗	PetscCall(PetscLogCpuToGpu(k*sizeof(PetscScalar)));
61		✗	}
62		✗	PetscCall(BVMultVec_BLAS_HIP(X,X->n,X->k-X->l,alpha,d_px+(X->nc+X->l)*X->ld,X->ld,d_q,beta,d_py));
63		✗	PetscCall(VecHIPRestoreArrayRead(x->v,&d_px));
64		✗	if (beta==(PetscScalar)0.0) PetscCall(VecHIPRestoreArrayWrite(y,&d_py));
65		✗	else PetscCall(VecHIPRestoreArray(y,&d_py));
66		✗	if (!q) PetscCall(VecHIPRestoreArray(X->buffer,&d_q));
67		✗	else PetscCallHIP(hipFree(d_q));
68		✗	PetscFunctionReturn(PETSC_SUCCESS);
69		✗	}
70
71		✗	PetscErrorCode BVMultInPlace_Svec_HIP(BV V,Mat Q,PetscInt s,PetscInt e)
72			{
73		✗	BV_SVEC *ctx = (BV_SVEC*)V->data;
74		✗	PetscScalar *d_pv;
75		✗	const PetscScalar *d_q;
76		✗	PetscInt ldq;
77
78			PetscFunctionBegin;
79		✗	if (s>=e || !V->n) PetscFunctionReturn(PETSC_SUCCESS);
80		✗	PetscCall(MatDenseGetLDA(Q,&ldq));
81		✗	PetscCall(VecHIPGetArray(ctx->v,&d_pv));
82		✗	PetscCall(BV_MatDenseHIPGetArrayRead(V,Q,&d_q));
83		✗	PetscCall(BVMultInPlace_BLAS_HIP(V,V->n,V->k-V->l,s-V->l,e-V->l,d_pv+(V->nc+V->l)*V->ld,V->ld,d_q+V->l*ldq+V->l,ldq,PETSC_FALSE));
84		✗	PetscCall(BV_MatDenseHIPRestoreArrayRead(V,Q,&d_q));
85		✗	PetscCall(VecHIPRestoreArray(ctx->v,&d_pv));
86		✗	PetscFunctionReturn(PETSC_SUCCESS);
87		✗	}
88
89		✗	PetscErrorCode BVMultInPlaceHermitianTranspose_Svec_HIP(BV V,Mat Q,PetscInt s,PetscInt e)
90			{
91		✗	BV_SVEC *ctx = (BV_SVEC*)V->data;
92		✗	PetscScalar *d_pv;
93		✗	const PetscScalar *d_q;
94		✗	PetscInt ldq;
95
96			PetscFunctionBegin;
97		✗	if (s>=e || !V->n) PetscFunctionReturn(PETSC_SUCCESS);
98		✗	PetscCall(MatDenseGetLDA(Q,&ldq));
99		✗	PetscCall(VecHIPGetArray(ctx->v,&d_pv));
100		✗	PetscCall(BV_MatDenseHIPGetArrayRead(V,Q,&d_q));
101		✗	PetscCall(BVMultInPlace_BLAS_HIP(V,V->n,V->k-V->l,s-V->l,e-V->l,d_pv+(V->nc+V->l)*V->ld,V->ld,d_q+V->l*ldq+V->l,ldq,PETSC_TRUE));
102		✗	PetscCall(BV_MatDenseHIPRestoreArrayRead(V,Q,&d_q));
103		✗	PetscCall(VecHIPRestoreArray(ctx->v,&d_pv));
104		✗	PetscFunctionReturn(PETSC_SUCCESS);
105		✗	}
106
107		✗	PetscErrorCode BVDot_Svec_HIP(BV X,BV Y,Mat M)
108			{
109		✗	BV_SVEC *x = (BV_SVEC*)X->data,*y = (BV_SVEC*)Y->data;
110		✗	const PetscScalar *d_px,*d_py;
111		✗	PetscScalar *pm;
112		✗	PetscInt ldm;
113
114			PetscFunctionBegin;
115		✗	PetscCall(MatDenseGetLDA(M,&ldm));
116		✗	PetscCall(VecHIPGetArrayRead(x->v,&d_px));
117		✗	PetscCall(VecHIPGetArrayRead(y->v,&d_py));
118		✗	PetscCall(MatDenseGetArrayWrite(M,&pm));
119		✗	PetscCall(BVDot_BLAS_HIP(X,Y->k-Y->l,X->k-X->l,X->n,d_py+(Y->nc+Y->l)*Y->ld,Y->ld,d_px+(X->nc+X->l)*X->ld,X->ld,pm+X->l*ldm+Y->l,ldm,x->mpi));
120		✗	PetscCall(MatDenseRestoreArrayWrite(M,&pm));
121		✗	PetscCall(VecHIPRestoreArrayRead(x->v,&d_px));
122		✗	PetscCall(VecHIPRestoreArrayRead(y->v,&d_py));
123		✗	PetscFunctionReturn(PETSC_SUCCESS);
124		✗	}
125
126		✗	PetscErrorCode BVDotVec_Svec_HIP(BV X,Vec y,PetscScalar *q)
127			{
128		✗	BV_SVEC *x = (BV_SVEC*)X->data;
129		✗	const PetscScalar *d_px,*d_py;
130		✗	Vec z = y;
131
132			PetscFunctionBegin;
133		✗	if (PetscUnlikely(X->matrix)) {
134		✗	PetscCall(BV_IPMatMult(X,y));
135		✗	z = X->Bx;
136		✗	}
137		✗	PetscCall(VecHIPGetArrayRead(x->v,&d_px));
138		✗	PetscCall(VecHIPGetArrayRead(z,&d_py));
139		✗	PetscCall(BVDotVec_BLAS_HIP(X,X->n,X->k-X->l,d_px+(X->nc+X->l)*X->ld,X->ld,d_py,q,x->mpi));
140		✗	PetscCall(VecHIPRestoreArrayRead(z,&d_py));
141		✗	PetscCall(VecHIPRestoreArrayRead(x->v,&d_px));
142		✗	PetscFunctionReturn(PETSC_SUCCESS);
143		✗	}
144
145		✗	PetscErrorCode BVDotVec_Local_Svec_HIP(BV X,Vec y,PetscScalar *m)
146			{
147		✗	BV_SVEC *x = (BV_SVEC*)X->data;
148		✗	const PetscScalar *d_px,*d_py;
149		✗	Vec z = y;
150
151			PetscFunctionBegin;
152		✗	if (PetscUnlikely(X->matrix)) {
153		✗	PetscCall(BV_IPMatMult(X,y));
154		✗	z = X->Bx;
155		✗	}
156		✗	PetscCall(VecHIPGetArrayRead(x->v,&d_px));
157		✗	PetscCall(VecHIPGetArrayRead(z,&d_py));
158		✗	PetscCall(BVDotVec_BLAS_HIP(X,X->n,X->k-X->l,d_px+(X->nc+X->l)*X->ld,X->ld,d_py,m,PETSC_FALSE));
159		✗	PetscCall(VecHIPRestoreArrayRead(z,&d_py));
160		✗	PetscCall(VecHIPRestoreArrayRead(x->v,&d_px));
161		✗	PetscFunctionReturn(PETSC_SUCCESS);
162		✗	}
163
164		✗	PetscErrorCode BVScale_Svec_HIP(BV bv,PetscInt j,PetscScalar alpha)
165			{
166		✗	BV_SVEC *ctx = (BV_SVEC*)bv->data;
167		✗	PetscScalar *d_array,*d_A;
168		✗	PetscInt n=0;
169
170			PetscFunctionBegin;
171		✗	if (!bv->n) PetscFunctionReturn(PETSC_SUCCESS);
172		✗	PetscCall(VecHIPGetArray(ctx->v,&d_array));
173		✗	if (PetscUnlikely(j<0)) {
174		✗	d_A = d_array+(bv->nc+bv->l)*bv->ld;
175		✗	n = (bv->k-bv->l)*bv->ld;
176		✗	} else {
177		✗	d_A = d_array+(bv->nc+j)*bv->ld;
178		✗	n = bv->n;
179			}
180		✗	PetscCall(BVScale_BLAS_HIP(bv,n,d_A,alpha));
181		✗	PetscCall(VecHIPRestoreArray(ctx->v,&d_array));
182		✗	PetscFunctionReturn(PETSC_SUCCESS);
183		✗	}
184
185		✗	PetscErrorCode BVNorm_Svec_HIP(BV bv,PetscInt j,NormType type,PetscReal *val)
186			{
187		✗	BV_SVEC *ctx = (BV_SVEC*)bv->data;
188		✗	const PetscScalar *array,*d_array,*d_A;
189		✗	PetscInt n=0;
190
191			PetscFunctionBegin;
192		✗	if (!ctx->mpi && ((j<0 && type==NORM_FROBENIUS && bv->ld==bv->n) || (j>=0 && type==NORM_2))) {
193			/* compute on GPU with hipBLAS - TODO: include the MPI case here */
194		✗	*val = 0.0;
195		✗	if (!bv->n) PetscFunctionReturn(PETSC_SUCCESS);
196		✗	PetscCall(VecHIPGetArrayRead(ctx->v,&d_array));
197		✗	if (PetscUnlikely(j<0)) {
198		✗	d_A = d_array+(bv->nc+bv->l)*bv->ld;
199		✗	n = (bv->k-bv->l)*bv->ld;
200		✗	} else {
201		✗	d_A = d_array+(bv->nc+j)*bv->ld;
202		✗	n = bv->n;
203			}
204		✗	PetscCall(BVNorm_BLAS_HIP(bv,n,d_A,val));
205		✗	PetscCall(VecHIPRestoreArrayRead(ctx->v,&d_array));
206		✗	} else {
207			/* compute on CPU */
208		✗	PetscCall(VecGetArrayRead(ctx->v,&array));
209		✗	if (PetscUnlikely(j<0)) PetscCall(BVNorm_LAPACK_Private(bv,bv->n,bv->k-bv->l,array+(bv->nc+bv->l)*bv->ld,bv->ld,type,val,ctx->mpi));
210		✗	else PetscCall(BVNorm_LAPACK_Private(bv,bv->n,1,array+(bv->nc+j)*bv->ld,bv->ld,type,val,ctx->mpi));
211		✗	PetscCall(VecRestoreArrayRead(ctx->v,&array));
212			}
213		✗	PetscFunctionReturn(PETSC_SUCCESS);
214		✗	}
215
216		✗	PetscErrorCode BVNorm_Local_Svec_HIP(BV bv,PetscInt j,NormType type,PetscReal *val)
217			{
218		✗	BV_SVEC *ctx = (BV_SVEC*)bv->data;
219		✗	const PetscScalar *array,*d_array,*d_A;
220		✗	PetscInt n=0;
221
222			PetscFunctionBegin;
223		✗	if ((j<0 && type==NORM_FROBENIUS && bv->ld==bv->n) || (j>=0 && type==NORM_2)) {
224			/* compute on GPU with hipBLAS */
225		✗	*val = 0.0;
226		✗	if (!bv->n) PetscFunctionReturn(PETSC_SUCCESS);
227		✗	PetscCall(VecHIPGetArrayRead(ctx->v,&d_array));
228		✗	if (PetscUnlikely(j<0)) {
229		✗	d_A = d_array+(bv->nc+bv->l)*bv->ld;
230		✗	n = (bv->k-bv->l)*bv->ld;
231		✗	} else {
232		✗	d_A = d_array+(bv->nc+j)*bv->ld;
233		✗	n = bv->n;
234			}
235		✗	PetscCall(BVNorm_BLAS_HIP(bv,n,d_A,val));
236		✗	PetscCall(VecHIPRestoreArrayRead(ctx->v,&d_array));
237		✗	} else {
238			/* compute on CPU */
239		✗	PetscCall(VecGetArrayRead(ctx->v,&array));
240		✗	if (PetscUnlikely(j<0)) PetscCall(BVNorm_LAPACK_Private(bv,bv->n,bv->k-bv->l,array+(bv->nc+bv->l)*bv->ld,bv->ld,type,val,PETSC_FALSE));
241		✗	else PetscCall(BVNorm_LAPACK_Private(bv,bv->n,1,array+(bv->nc+j)*bv->ld,bv->ld,type,val,PETSC_FALSE));
242		✗	PetscCall(VecRestoreArrayRead(ctx->v,&array));
243			}
244		✗	PetscFunctionReturn(PETSC_SUCCESS);
245		✗	}
246
247		✗	PetscErrorCode BVNormalize_Svec_HIP(BV bv,PetscScalar *eigi)
248			{
249		✗	BV_SVEC *ctx = (BV_SVEC*)bv->data;
250		✗	PetscScalar *array,*d_array,*wi=NULL;
251
252			PetscFunctionBegin;
253		✗	if (eigi) wi = eigi+bv->l;
254		✗	if (!ctx->mpi) {
255			/* compute on GPU with hipBLAS - TODO: include the MPI case here */
256		✗	if (!bv->n) PetscFunctionReturn(PETSC_SUCCESS);
257		✗	PetscCall(VecHIPGetArray(ctx->v,&d_array));
258		✗	PetscCall(BVNormalize_BLAS_HIP(bv,bv->n,bv->k-bv->l,d_array+(bv->nc+bv->l)*bv->ld,bv->ld,wi));
259		✗	PetscCall(VecHIPRestoreArray(ctx->v,&d_array));
260		✗	} else {
261			/* compute on CPU */
262		✗	PetscCall(VecGetArray(ctx->v,&array));
263		✗	PetscCall(BVNormalize_LAPACK_Private(bv,bv->n,bv->k-bv->l,array+(bv->nc+bv->l)*bv->ld,bv->ld,wi,ctx->mpi));
264		✗	PetscCall(VecRestoreArray(ctx->v,&array));
265			}
266		✗	PetscFunctionReturn(PETSC_SUCCESS);
267		✗	}
268
269		✗	PetscErrorCode BVMatMult_Svec_HIP(BV V,Mat A,BV W)
270			{
271		✗	BV_SVEC *v = (BV_SVEC*)V->data,*w = (BV_SVEC*)W->data;
272		✗	Mat Vmat,Wmat;
273		✗	const PetscScalar *d_pv;
274		✗	PetscScalar *d_pw;
275		✗	PetscInt j;
276
277			PetscFunctionBegin;
278		✗	if (V->vmm) {
279		✗	PetscCall(BVGetMat(V,&Vmat));
280		✗	PetscCall(BVGetMat(W,&Wmat));
281		✗	PetscCall(MatProductCreateWithMat(A,Vmat,NULL,Wmat));
282		✗	PetscCall(MatProductSetType(Wmat,MATPRODUCT_AB));
283		✗	PetscCall(MatProductSetFromOptions(Wmat));
284		✗	PetscCall(MatProductSymbolic(Wmat));
285		✗	PetscCall(MatProductNumeric(Wmat));
286		✗	PetscCall(MatProductClear(Wmat));
287		✗	PetscCall(BVRestoreMat(V,&Vmat));
288		✗	PetscCall(BVRestoreMat(W,&Wmat));
289		✗	} else {
290		✗	PetscCall(VecHIPGetArrayRead(v->v,&d_pv));
291		✗	PetscCall(VecHIPGetArrayWrite(w->v,&d_pw));
292		✗	for (j=0;j<V->k-V->l;j++) {
293		✗	PetscCall(VecHIPPlaceArray(V->cv[1],(PetscScalar *)d_pv+(V->nc+V->l+j)*V->ld));
294		✗	PetscCall(VecHIPPlaceArray(W->cv[1],d_pw+(W->nc+W->l+j)*W->ld));
295		✗	PetscCall(MatMult(A,V->cv[1],W->cv[1]));
296		✗	PetscCall(VecHIPResetArray(V->cv[1]));
297		✗	PetscCall(VecHIPResetArray(W->cv[1]));
298		✗	}
299		✗	PetscCall(VecHIPRestoreArrayRead(v->v,&d_pv));
300		✗	PetscCall(VecHIPRestoreArrayWrite(w->v,&d_pw));
301			}
302		✗	PetscFunctionReturn(PETSC_SUCCESS);
303		✗	}
304
305		✗	PetscErrorCode BVCopy_Svec_HIP(BV V,BV W)
306			{
307		✗	BV_SVEC *v = (BV_SVEC*)V->data,*w = (BV_SVEC*)W->data;
308		✗	const PetscScalar *d_pv;
309		✗	PetscScalar *d_pw;
310
311			PetscFunctionBegin;
312		✗	PetscCall(VecHIPGetArrayRead(v->v,&d_pv));
313		✗	PetscCall(VecHIPGetArray(w->v,&d_pw));
314		✗	PetscCallHIP(hipMemcpy2D(d_pw+(W->nc+W->l)*W->ld,W->ld*sizeof(PetscScalar),d_pv+(V->nc+V->l)*V->ld,V->ld*sizeof(PetscScalar),V->n*sizeof(PetscScalar),V->k-V->l,hipMemcpyDeviceToDevice));
315		✗	PetscCall(VecHIPRestoreArrayRead(v->v,&d_pv));
316		✗	PetscCall(VecHIPRestoreArray(w->v,&d_pw));
317		✗	PetscFunctionReturn(PETSC_SUCCESS);
318		✗	}
319
320		✗	PetscErrorCode BVCopyColumn_Svec_HIP(BV V,PetscInt j,PetscInt i)
321			{
322		✗	BV_SVEC *v = (BV_SVEC*)V->data;
323		✗	PetscScalar *d_pv;
324
325			PetscFunctionBegin;
326		✗	PetscCall(VecHIPGetArray(v->v,&d_pv));
327		✗	PetscCallHIP(hipMemcpy(d_pv+(V->nc+i)*V->ld,d_pv+(V->nc+j)*V->ld,V->n*sizeof(PetscScalar),hipMemcpyDeviceToDevice));
328		✗	PetscCall(VecHIPRestoreArray(v->v,&d_pv));
329		✗	PetscFunctionReturn(PETSC_SUCCESS);
330		✗	}
331
332		✗	PetscErrorCode BVResize_Svec_HIP(BV bv,PetscInt m,PetscBool copy)
333			{
334		✗	BV_SVEC *ctx = (BV_SVEC*)bv->data;
335		✗	const PetscScalar *d_pv;
336		✗	PetscScalar *d_pnew;
337		✗	PetscInt bs;
338		✗	Vec vnew;
339		✗	char str[50];
340
341			PetscFunctionBegin;
342		✗	PetscCall(PetscLayoutGetBlockSize(bv->map,&bs));
343		✗	PetscCall(VecCreate(PetscObjectComm((PetscObject)bv),&vnew));
344		✗	PetscCall(VecSetType(vnew,bv->vtype));
345		✗	PetscCall(VecSetSizes(vnew,m*bv->ld,PETSC_DECIDE));
346		✗	PetscCall(VecSetBlockSize(vnew,bs));
347		✗	if (((PetscObject)bv)->name) {
348		✗	PetscCall(PetscSNPrintf(str,sizeof(str),"%s_0",((PetscObject)bv)->name));
349		✗	PetscCall(PetscObjectSetName((PetscObject)vnew,str));
350		✗	}
351		✗	if (copy) {
352		✗	PetscCall(VecHIPGetArrayRead(ctx->v,&d_pv));
353		✗	PetscCall(VecHIPGetArrayWrite(vnew,&d_pnew));
354		✗	PetscCallHIP(hipMemcpy(d_pnew,d_pv,PetscMin(m,bv->m)*bv->ld*sizeof(PetscScalar),hipMemcpyDeviceToDevice));
355		✗	PetscCall(VecHIPRestoreArrayRead(ctx->v,&d_pv));
356		✗	PetscCall(VecHIPRestoreArrayWrite(vnew,&d_pnew));
357		✗	}
358		✗	PetscCall(VecDestroy(&ctx->v));
359		✗	ctx->v = vnew;
360		✗	PetscFunctionReturn(PETSC_SUCCESS);
361		✗	}
362
363		✗	PetscErrorCode BVGetColumn_Svec_HIP(BV bv,PetscInt j,Vec*)
364			{
365		✗	BV_SVEC *ctx = (BV_SVEC*)bv->data;
366		✗	PetscScalar *d_pv;
367		✗	PetscInt l;
368
369			PetscFunctionBegin;
370		✗	l = BVAvailableVec;
371		✗	PetscCall(VecHIPGetArray(ctx->v,&d_pv));
372		✗	PetscCall(VecHIPPlaceArray(bv->cv[l],d_pv+(bv->nc+j)*bv->ld));
373		✗	PetscFunctionReturn(PETSC_SUCCESS);
374		✗	}
375
376		✗	PetscErrorCode BVRestoreColumn_Svec_HIP(BV bv,PetscInt j,Vec*)
377			{
378		✗	BV_SVEC *ctx = (BV_SVEC*)bv->data;
379		✗	PetscInt l;
380
381			PetscFunctionBegin;
382		✗	l = (j==bv->ci[0])? 0: 1;
383		✗	PetscCall(VecHIPResetArray(bv->cv[l]));
384		✗	PetscCall(VecHIPRestoreArray(ctx->v,NULL));
385		✗	PetscFunctionReturn(PETSC_SUCCESS);
386		✗	}
387
388		✗	PetscErrorCode BVRestoreSplit_Svec_HIP(BV bv,BV *L,BV *R)
389			{
390		✗	Vec v;
391		✗	const PetscScalar *d_pv;
392		✗	PetscObjectState lstate,rstate;
393		✗	PetscBool change=PETSC_FALSE;
394
395			PetscFunctionBegin;
396			/* force sync flag to PETSC_OFFLOAD_BOTH */
397		✗	if (L) {
398		✗	PetscCall(PetscObjectStateGet((PetscObject)*L,&lstate));
399		✗	if (lstate != bv->lstate) {
400		✗	v = ((BV_SVEC*)bv->L->data)->v;
401		✗	PetscCall(VecHIPGetArrayRead(v,&d_pv));
402		✗	PetscCall(VecHIPRestoreArrayRead(v,&d_pv));
403		✗	change = PETSC_TRUE;
404		✗	}
405		✗	}
406		✗	if (R) {
407		✗	PetscCall(PetscObjectStateGet((PetscObject)*R,&rstate));
408		✗	if (rstate != bv->rstate) {
409		✗	v = ((BV_SVEC*)bv->R->data)->v;
410		✗	PetscCall(VecHIPGetArrayRead(v,&d_pv));
411		✗	PetscCall(VecHIPRestoreArrayRead(v,&d_pv));
412		✗	change = PETSC_TRUE;
413		✗	}
414		✗	}
415		✗	if (change) {
416		✗	v = ((BV_SVEC*)bv->data)->v;
417		✗	PetscCall(VecHIPGetArray(v,(PetscScalar **)&d_pv));
418		✗	PetscCall(VecHIPRestoreArray(v,(PetscScalar **)&d_pv));
419		✗	}
420		✗	PetscFunctionReturn(PETSC_SUCCESS);
421		✗	}
422
423		✗	PetscErrorCode BVRestoreSplitRows_Svec_HIP(BV bv,IS,IS,BV *U,BV *L)
424			{
425		✗	Vec v;
426		✗	const PetscScalar *d_pv;
427		✗	PetscObjectState lstate,rstate;
428		✗	PetscBool change=PETSC_FALSE;
429
430			PetscFunctionBegin;
431			/* force sync flag to PETSC_OFFLOAD_BOTH */
432		✗	if (U) {
433		✗	PetscCall(PetscObjectStateGet((PetscObject)*U,&rstate));
434		✗	if (rstate != bv->rstate) {
435		✗	v = ((BV_SVEC*)bv->R->data)->v;
436		✗	PetscCall(VecHIPGetArrayRead(v,&d_pv));
437		✗	PetscCall(VecHIPRestoreArrayRead(v,&d_pv));
438		✗	change = PETSC_TRUE;
439		✗	}
440		✗	}
441		✗	if (L) {
442		✗	PetscCall(PetscObjectStateGet((PetscObject)*L,&lstate));
443		✗	if (lstate != bv->lstate) {
444		✗	v = ((BV_SVEC*)bv->L->data)->v;
445		✗	PetscCall(VecHIPGetArrayRead(v,&d_pv));
446		✗	PetscCall(VecHIPRestoreArrayRead(v,&d_pv));
447		✗	change = PETSC_TRUE;
448		✗	}
449		✗	}
450		✗	if (change) {
451		✗	v = ((BV_SVEC*)bv->data)->v;
452		✗	PetscCall(VecHIPGetArray(v,(PetscScalar **)&d_pv));
453		✗	PetscCall(VecHIPRestoreArray(v,(PetscScalar **)&d_pv));
454		✗	}
455		✗	PetscFunctionReturn(PETSC_SUCCESS);
456		✗	}
457
458		✗	PetscErrorCode BVGetMat_Svec_HIP(BV bv,Mat *A)
459			{
460		✗	BV_SVEC *ctx = (BV_SVEC*)bv->data;
461		✗	PetscScalar *vv,*aa;
462		✗	PetscBool create=PETSC_FALSE;
463		✗	PetscInt m,cols;
464
465			PetscFunctionBegin;
466		✗	m = bv->k-bv->l;
467		✗	if (!bv->Aget) create=PETSC_TRUE;
468			else {
469		✗	PetscCall(MatDenseHIPGetArray(bv->Aget,&aa));
470		✗	PetscCheck(!aa,PetscObjectComm((PetscObject)bv),PETSC_ERR_ARG_WRONGSTATE,"BVGetMat already called on this BV");
471		✗	PetscCall(MatGetSize(bv->Aget,NULL,&cols));
472		✗	if (cols!=m) {
473		✗	PetscCall(MatDestroy(&bv->Aget));
474		✗	create=PETSC_TRUE;
475		✗	}
476			}
477		✗	PetscCall(VecHIPGetArray(ctx->v,&vv));
478		✗	if (create) {
479		✗	PetscCall(MatCreateDenseFromVecType(PetscObjectComm((PetscObject)bv),bv->vtype,bv->n,PETSC_DECIDE,bv->N,m,bv->ld,vv,&bv->Aget)); /* pass a pointer to avoid allocation of storage */
480		✗	PetscCall(MatDenseHIPReplaceArray(bv->Aget,NULL)); /* replace with a null pointer, the value after BVRestoreMat */
481		✗	}
482		✗	PetscCall(MatDenseHIPPlaceArray(bv->Aget,vv+(bv->nc+bv->l)*bv->ld)); /* set the actual pointer */
483		✗	*A = bv->Aget;
484		✗	PetscFunctionReturn(PETSC_SUCCESS);
485		✗	}
486
487		✗	PetscErrorCode BVRestoreMat_Svec_HIP(BV bv,Mat *A)
488			{
489		✗	BV_SVEC *ctx = (BV_SVEC*)bv->data;
490		✗	PetscScalar *vv,*aa;
491
492			PetscFunctionBegin;
493		✗	PetscCall(MatDenseHIPGetArray(bv->Aget,&aa));
494		✗	vv = aa-(bv->nc+bv->l)*bv->ld;
495		✗	PetscCall(MatDenseHIPResetArray(bv->Aget));
496		✗	PetscCall(VecHIPRestoreArray(ctx->v,&vv));
497		✗	*A = NULL;
498		✗	PetscFunctionReturn(PETSC_SUCCESS);
499		✗	}
500