327 …srcranks, PetscMPIInt *srcranks, const char *src, PetscInt *srcdisp, char *dst, PetscInt *dstdisp,…  in GetDataFromRemotelyAccessible()  argument
334 …nvshmem_getmem_nbi(dst + (dstdisp[bid] - dstdisp[0]) * unitbytes, src + srcdisp[bid] * unitbytes, …  in GetDataFromRemotelyAccessible()
346   char        *src, *dst;  in PetscSFLinkGetDataBegin_NVSHMEM()  local
366 …dst       = link->leafbuf[PETSCSF_REMOTE][PETSC_MEMTYPE_DEVICE]; /* recv buf is the local leaf buf…  in PetscSFLinkGetDataBegin_NVSHMEM()
384 …dst       = link->rootbuf[PETSCSF_REMOTE][PETSC_MEMTYPE_DEVICE]; /* the local root buf is the recv…  in PetscSFLinkGetDataBegin_NVSHMEM()
415 …ks, 1, 0, link->remoteCommStream>>>(nsrcranks, srcranks_d, src, srcdisp_d, dst, dstdisp_d, link->u…  in PetscSFLinkGetDataBegin_NVSHMEM()
425 …nvshmemx_getmem_nbi_on_stream(dst + (dstdisp_h[i] - dstdisp_h[0]) * link->unitbytes, src + srcdisp…  in PetscSFLinkGetDataBegin_NVSHMEM()
488 …tDataToRemotelyAccessible(PetscInt ndstranks, PetscMPIInt *dstranks, char *dst, PetscInt *dstdisp,…  in WaitAndPutDataToRemotelyAccessible()  argument
493   if (!nvshmem_ptr(dst, pe)) {  in WaitAndPutDataToRemotelyAccessible()
497 …nvshmem_putmem_nbi(dst + dstdisp[bid] * unitbytes, src + (srcdisp[bid] - srcdisp[0]) * unitbytes, …  in WaitAndPutDataToRemotelyAccessible()
502 …FromLocallyAccessible(PetscInt ndstranks, PetscMPIInt *dstranks, uint64_t *srcsig, const char *dst)  in WaitSignalsFromLocallyAccessible()  argument
506     if (nvshmem_ptr(dst, pe)) {  in WaitSignalsFromLocallyAccessible()
519   char          *src, *dst;  in PetscSFLinkPutDataBegin_NVSHMEM()  local
531     dst       = link->leafbuf[PETSCSF_REMOTE][PETSC_MEMTYPE_DEVICE];  in PetscSFLinkPutDataBegin_NVSHMEM()
544     dst       = link->rootbuf[PETSCSF_REMOTE][PETSC_MEMTYPE_DEVICE];  in PetscSFLinkPutDataBegin_NVSHMEM()
560     if (nvshmem_ptr(dst, dstranks_h[i])) nLocallyAccessible++;  in PetscSFLinkPutDataBegin_NVSHMEM()
565 …ssible<<<ndstranks, 1, 0, link->remoteCommStream>>>(ndstranks, dstranks_d, dst, dstdisp_d, src, sr…  in PetscSFLinkPutDataBegin_NVSHMEM()
571 …nalsFromLocallyAccessible<<<1, 1, 0, link->remoteCommStream>>>(ndstranks, dstranks_d, srcsig, dst);  in PetscSFLinkPutDataBegin_NVSHMEM()
574 …    if (nvshmem_ptr(dst, pe)) { /* If return a non-null pointer, then <pe> is locally accessible */  in PetscSFLinkPutDataBegin_NVSHMEM()
577 …nvshmemx_putmem_nbi_on_stream(dst + dstdisp_h[i] * link->unitbytes, src + (srcdisp_h[i] - srcdisp_…  in PetscSFLinkPutDataBegin_NVSHMEM()
793 PetscErrorCode PetscNvshmemSum(PetscInt count, float *dst, const float *src)  in PetscNvshmemSum()  argument
799   nvshmemx_float_sum_reduce_on_stream(NVSHMEM_TEAM_WORLD, dst, src, num, PetscDefaultCudaStream);  in PetscNvshmemSum()
803 PetscErrorCode PetscNvshmemMax(PetscInt count, float *dst, const float *src)  in PetscNvshmemMax()  argument
809   nvshmemx_float_max_reduce_on_stream(NVSHMEM_TEAM_WORLD, dst, src, num, PetscDefaultCudaStream);  in PetscNvshmemMax()
813 PetscErrorCode PetscNvshmemSum(PetscInt count, double *dst, const double *src)  in PetscNvshmemSum()  argument
819   nvshmemx_double_sum_reduce_on_stream(NVSHMEM_TEAM_WORLD, dst, src, num, PetscDefaultCudaStream);  in PetscNvshmemSum()
823 PetscErrorCode PetscNvshmemMax(PetscInt count, double *dst, const double *src)  in PetscNvshmemMax()  argument
829   nvshmemx_double_max_reduce_on_stream(NVSHMEM_TEAM_WORLD, dst, src, num, PetscDefaultCudaStream);  in PetscNvshmemMax()