Actual source code: mpihip.hip.cpp


  2: /*
  3:    This file contains routines for Parallel vector operations.
  4:  */
  5: #define PETSC_SKIP_SPINLOCK

  7: #include <petscconf.h>
  8: #include <../src/vec/vec/impls/mpi/pvecimpl.h>
  9: #include <petsc/private/hipvecimpl.h>

 11: /*MC
 12:    VECHIP - VECHIP = "hip" - A VECSEQHIP on a single-process communicator, and VECMPIHIP otherwise.

 14:    Options Database Keys:
 15: . -vec_type hip - sets the vector type to VECHIP during a call to VecSetFromOptions()

 17:   Level: beginner

 19: .seealso: VecCreate(), VecSetType(), VecSetFromOptions(), VecCreateMPIWithArray(), VECSEQHIP, VECMPIHIP, VECSTANDARD, VecType, VecCreateMPI(), VecSetPinnedMemoryMin()
 20: M*/

 22: PetscErrorCode VecDestroy_MPIHIP(Vec v)
 23: {
 24:   Vec_MPI        *vecmpi = (Vec_MPI*)v->data;
 25:   Vec_HIP        *vechip;

 27:   if (v->spptr) {
 28:     vechip = (Vec_HIP*)v->spptr;
 29:     if (vechip->GPUarray_allocated) {
 30:       hipFree(vechip->GPUarray_allocated);
 31:       vechip->GPUarray_allocated = NULL;
 32:     }
 33:     if (vechip->stream) {
 34:       hipStreamDestroy(vechip->stream);
 35:     }
 36:     if (v->pinned_memory) {
 37:       PetscMallocSetHIPHost();
 38:       PetscFree(vecmpi->array_allocated);
 39:       PetscMallocResetHIPHost();
 40:       v->pinned_memory = PETSC_FALSE;
 41:     }
 42:     PetscFree(v->spptr);
 43:   }
 44:   VecDestroy_MPI(v);
 45:   return 0;
 46: }

 48: PetscErrorCode VecNorm_MPIHIP(Vec xin,NormType type,PetscReal *z)
 49: {
 50:   PetscReal      sum,work = 0.0;

 52:   if (type == NORM_2 || type == NORM_FROBENIUS) {
 53:     VecNorm_SeqHIP(xin,NORM_2,&work);
 54:     work *= work;
 55:     MPIU_Allreduce(&work,&sum,1,MPIU_REAL,MPIU_SUM,PetscObjectComm((PetscObject)xin));
 56:     *z    = PetscSqrtReal(sum);
 57:   } else if (type == NORM_1) {
 58:     /* Find the local part */
 59:     VecNorm_SeqHIP(xin,NORM_1,&work);
 60:     /* Find the global max */
 61:     MPIU_Allreduce(&work,z,1,MPIU_REAL,MPIU_SUM,PetscObjectComm((PetscObject)xin));
 62:   } else if (type == NORM_INFINITY) {
 63:     /* Find the local max */
 64:     VecNorm_SeqHIP(xin,NORM_INFINITY,&work);
 65:     /* Find the global max */
 66:     MPIU_Allreduce(&work,z,1,MPIU_REAL,MPIU_MAX,PetscObjectComm((PetscObject)xin));
 67:   } else if (type == NORM_1_AND_2) {
 68:     PetscReal temp[2];
 69:     VecNorm_SeqHIP(xin,NORM_1,temp);
 70:     VecNorm_SeqHIP(xin,NORM_2,temp+1);
 71:     temp[1] = temp[1]*temp[1];
 72:     MPIU_Allreduce(temp,z,2,MPIU_REAL,MPIU_SUM,PetscObjectComm((PetscObject)xin));
 73:     z[1] = PetscSqrtReal(z[1]);
 74:   }
 75:   return 0;
 76: }

 78: PetscErrorCode VecDot_MPIHIP(Vec xin,Vec yin,PetscScalar *z)
 79: {
 80:   PetscScalar    sum,work;

 82:   VecDot_SeqHIP(xin,yin,&work);
 83:   MPIU_Allreduce(&work,&sum,1,MPIU_SCALAR,MPIU_SUM,PetscObjectComm((PetscObject)xin));
 84:   *z   = sum;
 85:   return 0;
 86: }

 88: PetscErrorCode VecTDot_MPIHIP(Vec xin,Vec yin,PetscScalar *z)
 89: {
 90:   PetscScalar    sum,work;

 92:   VecTDot_SeqHIP(xin,yin,&work);
 93:   MPIU_Allreduce(&work,&sum,1,MPIU_SCALAR,MPIU_SUM,PetscObjectComm((PetscObject)xin));
 94:   *z   = sum;
 95:   return 0;
 96: }

 98: PetscErrorCode VecMDot_MPIHIP(Vec xin,PetscInt nv,const Vec y[],PetscScalar *z)
 99: {
100:   PetscScalar    awork[128],*work = awork;

102:   if (nv > 128) {
103:     PetscMalloc1(nv,&work);
104:   }
105:   VecMDot_SeqHIP(xin,nv,y,work);
106:   MPIU_Allreduce(work,z,nv,MPIU_SCALAR,MPIU_SUM,PetscObjectComm((PetscObject)xin));
107:   if (nv > 128) {
108:     PetscFree(work);
109:   }
110:   return 0;
111: }

113: /*MC
114:    VECMPIHIP - VECMPIHIP = "mpihip" - The basic parallel vector, modified to use HIP

116:    Options Database Keys:
117: . -vec_type mpihip - sets the vector type to VECMPIHIP during a call to VecSetFromOptions()

119:   Level: beginner

121: .seealso: VecCreate(), VecSetType(), VecSetFromOptions(), VecCreateMPIWithArray(), VECMPI, VecType, VecCreateMPI(), VecSetPinnedMemoryMin()
122: M*/

124: PetscErrorCode VecDuplicate_MPIHIP(Vec win,Vec *v)
125: {
126:   Vec_MPI        *vw,*w = (Vec_MPI*)win->data;
127:   PetscScalar    *array;

129:   VecCreate(PetscObjectComm((PetscObject)win),v);
130:   PetscLayoutReference(win->map,&(*v)->map);

132:   VecCreate_MPIHIP_Private(*v,PETSC_TRUE,w->nghost,0);
133:   vw   = (Vec_MPI*)(*v)->data;
134:   PetscMemcpy((*v)->ops,win->ops,sizeof(struct _VecOps));

136:   /* save local representation of the parallel vector (and scatter) if it exists */
137:   if (w->localrep) {
138:     VecGetArray(*v,&array);
139:     VecCreateSeqWithArray(PETSC_COMM_SELF,1,win->map->n+w->nghost,array,&vw->localrep);
140:     PetscMemcpy(vw->localrep->ops,w->localrep->ops,sizeof(struct _VecOps));
141:     VecRestoreArray(*v,&array);
142:     PetscLogObjectParent((PetscObject)*v,(PetscObject)vw->localrep);
143:     vw->localupdate = w->localupdate;
144:     if (vw->localupdate) {
145:       PetscObjectReference((PetscObject)vw->localupdate);
146:     }
147:   }

149:   /* New vector should inherit stashing property of parent */
150:   (*v)->stash.donotstash   = win->stash.donotstash;
151:   (*v)->stash.ignorenegidx = win->stash.ignorenegidx;

153:   /* change type_name appropriately */
154:   VecHIPAllocateCheck(*v);
155:   PetscObjectChangeTypeName((PetscObject)(*v),VECMPIHIP);

157:   PetscObjectListDuplicate(((PetscObject)win)->olist,&((PetscObject)(*v))->olist);
158:   PetscFunctionListDuplicate(((PetscObject)win)->qlist,&((PetscObject)(*v))->qlist);
159:   (*v)->map->bs   = PetscAbs(win->map->bs);
160:   (*v)->bstash.bs = win->bstash.bs;
161:   return 0;
162: }

164: PetscErrorCode VecDotNorm2_MPIHIP(Vec s,Vec t,PetscScalar *dp,PetscScalar *nm)
165: {
166:   PetscScalar    work[2],sum[2];

168:   VecDotNorm2_SeqHIP(s,t,work,work+1);
169:   MPIU_Allreduce(&work,&sum,2,MPIU_SCALAR,MPIU_SUM,PetscObjectComm((PetscObject)s));
170:   *dp  = sum[0];
171:   *nm  = sum[1];
172:   return 0;
173: }

175: PetscErrorCode VecCreate_MPIHIP(Vec vv)
176: {
177:   PetscDeviceInitialize(PETSC_DEVICE_HIP);
178:   PetscLayoutSetUp(vv->map);
179:   VecHIPAllocateCheck(vv);
180:   VecCreate_MPIHIP_Private(vv,PETSC_FALSE,0,((Vec_HIP*)vv->spptr)->GPUarray_allocated);
181:   VecHIPAllocateCheckHost(vv);
182:   VecSet(vv,0.0);
183:   VecSet_Seq(vv,0.0);
184:   vv->offloadmask = PETSC_OFFLOAD_BOTH;
185:   return 0;
186: }

188: PetscErrorCode VecCreate_HIP(Vec v)
189: {
190:   PetscMPIInt    size;

192:   MPI_Comm_size(PetscObjectComm((PetscObject)v),&size);
193:   if (size == 1) {
194:     VecSetType(v,VECSEQHIP);
195:   } else {
196:     VecSetType(v,VECMPIHIP);
197:   }
198:   return 0;
199: }

201: /*@
202:  VecCreateMPIHIP - Creates a standard, parallel array-style vector for HIP devices.

204:  Collective

206:  Input Parameters:
207:  +  comm - the MPI communicator to use
208:  .  n - local vector length (or PETSC_DECIDE to have calculated if N is given)
209:  -  N - global vector length (or PETSC_DETERMINE to have calculated if n is given)

211:     Output Parameter:
212:  .  v - the vector

214:     Notes:
215:     Use VecDuplicate() or VecDuplicateVecs() to form additional vectors of the
216:     same type as an existing vector.

218:     Level: intermediate

220:  .seealso: VecCreateMPIHIPWithArray(), VecCreateMPIHIPWithArrays(), VecCreateSeqHIP(), VecCreateSeq(),
221:            VecCreateMPI(), VecCreate(), VecDuplicate(), VecDuplicateVecs(), VecCreateGhost(),
222:            VecCreateMPIWithArray(), VecCreateGhostWithArray(), VecMPISetGhost()

224:  @*/
225:  PetscErrorCode VecCreateMPIHIP(MPI_Comm comm,PetscInt n,PetscInt N,Vec *v)
226:  {
227:    VecCreate(comm,v);
228:    VecSetSizes(*v,n,N);
229:    VecSetType(*v,VECMPIHIP);
230:    return 0;
231:  }

233: /*@C
234:    VecCreateMPIHIPWithArray - Creates a parallel, array-style vector,
235:    where the user provides the GPU array space to store the vector values.

237:    Collective

239:    Input Parameters:
240: +  comm  - the MPI communicator to use
241: .  bs    - block size, same meaning as VecSetBlockSize()
242: .  n     - local vector length, cannot be PETSC_DECIDE
243: .  N     - global vector length (or PETSC_DECIDE to have calculated)
244: -  array - the user provided GPU array to store the vector values

246:    Output Parameter:
247: .  vv - the vector

249:    Notes:
250:    Use VecDuplicate() or VecDuplicateVecs() to form additional vectors of the
251:    same type as an existing vector.

253:    If the user-provided array is NULL, then VecHIPPlaceArray() can be used
254:    at a later stage to SET the array for storing the vector values.

256:    PETSc does NOT free the array when the vector is destroyed via VecDestroy().
257:    The user should not free the array until the vector is destroyed.

259:    Level: intermediate

261: .seealso: VecCreateSeqHIPWithArray(), VecCreateMPIWithArray(), VecCreateSeqWithArray(),
262:           VecCreate(), VecDuplicate(), VecDuplicateVecs(), VecCreateGhost(),
263:           VecCreateMPI(), VecCreateGhostWithArray(), VecPlaceArray()

265: @*/
266: PetscErrorCode  VecCreateMPIHIPWithArray(MPI_Comm comm,PetscInt bs,PetscInt n,PetscInt N,const PetscScalar array[],Vec *vv)
267: {
269:   PetscDeviceInitialize(PETSC_DEVICE_HIP);
270:   VecCreate(comm,vv);
271:   VecSetSizes(*vv,n,N);
272:   VecSetBlockSize(*vv,bs);
273:   VecCreate_MPIHIP_Private(*vv,PETSC_FALSE,0,array);
274:   return 0;
275: }

277: /*@C
278:    VecCreateMPIHIPWithArrays - Creates a parallel, array-style vector,
279:    where the user provides the GPU array space to store the vector values.

281:    Collective

283:    Input Parameters:
284: +  comm  - the MPI communicator to use
285: .  bs    - block size, same meaning as VecSetBlockSize()
286: .  n     - local vector length, cannot be PETSC_DECIDE
287: .  N     - global vector length (or PETSC_DECIDE to have calculated)
288: -  cpuarray - the user provided CPU array to store the vector values
289: -  gpuarray - the user provided GPU array to store the vector values

291:    Output Parameter:
292: .  vv - the vector

294:    Notes:
295:    If both cpuarray and gpuarray are provided, the caller must ensure that
296:    the provided arrays have identical values.

298:    Use VecDuplicate() or VecDuplicateVecs() to form additional vectors of the
299:    same type as an existing vector.

301:    PETSc does NOT free the provided arrays when the vector is destroyed via
302:    VecDestroy(). The user should not free the array until the vector is
303:    destroyed.

305:    Level: intermediate

307: .seealso: VecCreateSeqHIPWithArrays(), VecCreateMPIWithArray(), VecCreateSeqWithArray(),
308:           VecCreate(), VecDuplicate(), VecDuplicateVecs(), VecCreateGhost(),
309:           VecCreateMPI(), VecCreateGhostWithArray(), VecHIPPlaceArray(), VecPlaceArray(),
310:           VecHIPAllocateCheckHost()
311: @*/
312: PetscErrorCode  VecCreateMPIHIPWithArrays(MPI_Comm comm,PetscInt bs,PetscInt n,PetscInt N,const PetscScalar cpuarray[],const PetscScalar gpuarray[],Vec *vv)
313: {
314:   VecCreateMPIHIPWithArray(comm,bs,n,N,gpuarray,vv);

316:   if (cpuarray && gpuarray) {
317:     Vec_MPI *s         = (Vec_MPI*)((*vv)->data);
318:     s->array           = (PetscScalar*)cpuarray;
319:     (*vv)->offloadmask = PETSC_OFFLOAD_BOTH;
320:   } else if (cpuarray) {
321:     Vec_MPI *s         = (Vec_MPI*)((*vv)->data);
322:     s->array           = (PetscScalar*)cpuarray;
323:     (*vv)->offloadmask =  PETSC_OFFLOAD_CPU;
324:   } else if (gpuarray) {
325:     (*vv)->offloadmask = PETSC_OFFLOAD_GPU;
326:   } else {
327:     (*vv)->offloadmask = PETSC_OFFLOAD_UNALLOCATED;
328:   }

330:   return 0;
331: }

333: PetscErrorCode VecMax_MPIHIP(Vec xin,PetscInt *idx,PetscReal *z)
334: {
335:   PetscReal      work;

337:   VecMax_SeqHIP(xin,idx,&work);
338: #if defined(PETSC_HAVE_MPIUNI)
339:   *z = work;
340: #else
341:   if (!idx) {
342:     MPIU_Allreduce(&work,z,1,MPIU_REAL,MPIU_MAX,PetscObjectComm((PetscObject)xin));
343:   } else {
344:     struct { PetscReal v; PetscInt i; } in,out;

346:     in.v  = work;
347:     in.i  = *idx + xin->map->rstart;
348:     MPIU_Allreduce(&in,&out,1,MPIU_REAL_INT,MPIU_MAXLOC,PetscObjectComm((PetscObject)xin));
349:     *z    = out.v;
350:     *idx  = out.i;
351:   }
352: #endif
353:   return 0;
354: }

356: PetscErrorCode VecMin_MPIHIP(Vec xin,PetscInt *idx,PetscReal *z)
357: {
358:   PetscReal      work;

360:   VecMin_SeqHIP(xin,idx,&work);
361: #if defined(PETSC_HAVE_MPIUNI)
362:   *z = work;
363: #else
364:   if (!idx) {
365:     MPIU_Allreduce(&work,z,1,MPIU_REAL,MPIU_MIN,PetscObjectComm((PetscObject)xin));
366:   } else {
367:     struct { PetscReal v; PetscInt i; } in,out;

369:     in.v  = work;
370:     in.i  = *idx + xin->map->rstart;
371:     MPIU_Allreduce(&in,&out,1,MPIU_REAL_INT,MPIU_MINLOC,PetscObjectComm((PetscObject)xin));
372:     *z    = out.v;
373:     *idx  = out.i;
374:   }
375: #endif
376:   return 0;
377: }

379: PetscErrorCode VecBindToCPU_MPIHIP(Vec V,PetscBool bind)
380: {
381:   V->boundtocpu = bind;
382:   if (bind) {
383:     VecHIPCopyFromGPU(V);
384:     V->offloadmask = PETSC_OFFLOAD_CPU; /* since the CPU code will likely change values in the vector */
385:     V->ops->dotnorm2               = NULL;
386:     V->ops->waxpy                  = VecWAXPY_Seq;
387:     V->ops->dot                    = VecDot_MPI;
388:     V->ops->mdot                   = VecMDot_MPI;
389:     V->ops->tdot                   = VecTDot_MPI;
390:     V->ops->norm                   = VecNorm_MPI;
391:     V->ops->scale                  = VecScale_Seq;
392:     V->ops->copy                   = VecCopy_Seq;
393:     V->ops->set                    = VecSet_Seq;
394:     V->ops->swap                   = VecSwap_Seq;
395:     V->ops->axpy                   = VecAXPY_Seq;
396:     V->ops->axpby                  = VecAXPBY_Seq;
397:     V->ops->maxpy                  = VecMAXPY_Seq;
398:     V->ops->aypx                   = VecAYPX_Seq;
399:     V->ops->axpbypcz               = VecAXPBYPCZ_Seq;
400:     V->ops->pointwisemult          = VecPointwiseMult_Seq;
401:     V->ops->setrandom              = VecSetRandom_Seq;
402:     V->ops->placearray             = VecPlaceArray_Seq;
403:     V->ops->replacearray           = VecReplaceArray_SeqHIP;
404:     V->ops->resetarray             = VecResetArray_Seq;
405:     V->ops->dot_local              = VecDot_Seq;
406:     V->ops->tdot_local             = VecTDot_Seq;
407:     V->ops->norm_local             = VecNorm_Seq;
408:     V->ops->mdot_local             = VecMDot_Seq;
409:     V->ops->pointwisedivide        = VecPointwiseDivide_Seq;
410:     V->ops->getlocalvector         = NULL;
411:     V->ops->restorelocalvector     = NULL;
412:     V->ops->getlocalvectorread     = NULL;
413:     V->ops->restorelocalvectorread = NULL;
414:     V->ops->getarraywrite          = NULL;
415:     V->ops->getarrayandmemtype     = NULL;
416:     V->ops->restorearrayandmemtype = NULL;
417:     V->ops->getarraywriteandmemtype= NULL;
418:     V->ops->max                    = VecMax_MPI;
419:     V->ops->min                    = VecMin_MPI;
420:     V->ops->reciprocal             = VecReciprocal_Default;
421:     V->ops->sum                    = NULL;
422:     V->ops->shift                  = NULL;
423:   } else {
424:     V->ops->dotnorm2               = VecDotNorm2_MPIHIP;
425:     V->ops->waxpy                  = VecWAXPY_SeqHIP;
426:     V->ops->duplicate              = VecDuplicate_MPIHIP;
427:     V->ops->dot                    = VecDot_MPIHIP;
428:     V->ops->mdot                   = VecMDot_MPIHIP;
429:     V->ops->tdot                   = VecTDot_MPIHIP;
430:     V->ops->norm                   = VecNorm_MPIHIP;
431:     V->ops->scale                  = VecScale_SeqHIP;
432:     V->ops->copy                   = VecCopy_SeqHIP;
433:     V->ops->set                    = VecSet_SeqHIP;
434:     V->ops->swap                   = VecSwap_SeqHIP;
435:     V->ops->axpy                   = VecAXPY_SeqHIP;
436:     V->ops->axpby                  = VecAXPBY_SeqHIP;
437:     V->ops->maxpy                  = VecMAXPY_SeqHIP;
438:     V->ops->aypx                   = VecAYPX_SeqHIP;
439:     V->ops->axpbypcz               = VecAXPBYPCZ_SeqHIP;
440:     V->ops->pointwisemult          = VecPointwiseMult_SeqHIP;
441:     V->ops->setrandom              = VecSetRandom_SeqHIP;
442:     V->ops->placearray             = VecPlaceArray_SeqHIP;
443:     V->ops->replacearray           = VecReplaceArray_SeqHIP;
444:     V->ops->resetarray             = VecResetArray_SeqHIP;
445:     V->ops->dot_local              = VecDot_SeqHIP;
446:     V->ops->tdot_local             = VecTDot_SeqHIP;
447:     V->ops->norm_local             = VecNorm_SeqHIP;
448:     V->ops->mdot_local             = VecMDot_SeqHIP;
449:     V->ops->destroy                = VecDestroy_MPIHIP;
450:     V->ops->pointwisedivide        = VecPointwiseDivide_SeqHIP;
451:     V->ops->getlocalvector         = VecGetLocalVector_SeqHIP;
452:     V->ops->restorelocalvector     = VecRestoreLocalVector_SeqHIP;
453:     V->ops->getlocalvectorread     = VecGetLocalVectorRead_SeqHIP;
454:     V->ops->restorelocalvectorread = VecRestoreLocalVectorRead_SeqHIP;
455:     V->ops->getarraywrite          = VecGetArrayWrite_SeqHIP;
456:     V->ops->getarray               = VecGetArray_SeqHIP;
457:     V->ops->restorearray           = VecRestoreArray_SeqHIP;
458:     V->ops->getarrayandmemtype     = VecGetArrayAndMemType_SeqHIP;
459:     V->ops->restorearrayandmemtype = VecRestoreArrayAndMemType_SeqHIP;
460:     V->ops->getarraywriteandmemtype= VecGetArrayWriteAndMemType_SeqHIP;
461:     V->ops->max                    = VecMax_MPIHIP;
462:     V->ops->min                    = VecMin_MPIHIP;
463:     V->ops->reciprocal             = VecReciprocal_SeqHIP;
464:     V->ops->sum                    = VecSum_SeqHIP;
465:     V->ops->shift                  = VecShift_SeqHIP;
466:   }
467:   return 0;
468: }

470: PetscErrorCode VecCreate_MPIHIP_Private(Vec vv,PetscBool alloc,PetscInt nghost,const PetscScalar array[])
471: {
472:   Vec_HIP *vechip;

474:   VecCreate_MPI_Private(vv,PETSC_FALSE,0,0);
475:   PetscObjectChangeTypeName((PetscObject)vv,VECMPIHIP);

477:   VecBindToCPU_MPIHIP(vv,PETSC_FALSE);
478:   vv->ops->bindtocpu = VecBindToCPU_MPIHIP;

480:   /* Later, functions check for the Vec_HIP structure existence, so do not create it without array */
481:   if (alloc && !array) {
482:     VecHIPAllocateCheck(vv);
483:     VecHIPAllocateCheckHost(vv);
484:     VecSet(vv,0.0);
485:     VecSet_Seq(vv,0.0);
486:     vv->offloadmask = PETSC_OFFLOAD_BOTH;
487:   }
488:   if (array) {
489:     if (!vv->spptr) {
490:       PetscReal      pinned_memory_min;
491:       PetscBool      flag;

494:       /* Cannot use PetscNew() here because spptr is void* */
495:       PetscCalloc(sizeof(Vec_HIP),&vv->spptr);
496:       vechip = (Vec_HIP*)vv->spptr;
497:       vv->minimum_bytes_pinned_memory = 0;

499:       /* Need to parse command line for minimum size to use for pinned memory allocations on host here.
500:          Note: This same code duplicated in VecCreate_SeqHIP_Private() and VecHIPAllocateCheck(). Is there a good way to avoid this? */
501:       PetscOptionsBegin(PetscObjectComm((PetscObject)vv),((PetscObject)vv)->prefix,"VECHIP Options","Vec");
502:       pinned_memory_min = vv->minimum_bytes_pinned_memory;
503:       PetscOptionsReal("-vec_pinned_memory_min","Minimum size (in bytes) for an allocation to use pinned memory on host","VecSetPinnedMemoryMin",pinned_memory_min,&pinned_memory_min,&flag);
504:       if (flag) vv->minimum_bytes_pinned_memory = pinned_memory_min;
505:       PetscOptionsEnd();
506:     }
507:     vechip = (Vec_HIP*)vv->spptr;
508:     vechip->GPUarray = (PetscScalar*)array;
509:     vv->offloadmask = PETSC_OFFLOAD_GPU;
510:   }
511:   return 0;
512: }