2 #include <../src/vec/is/sf/impls/basic/sfpack.h>
327 …mRemotelyAccessible(PetscInt nsrcranks, PetscMPIInt *srcranks, const char *src, PetscInt *srcdisp,…  in GetDataFromRemotelyAccessible()  argument
332   if (!nvshmem_ptr(src, pe)) {  in GetDataFromRemotelyAccessible()
334 …nvshmem_getmem_nbi(dst + (dstdisp[bid] - dstdisp[0]) * unitbytes, src + srcdisp[bid] * unitbytes, …  in GetDataFromRemotelyAccessible()
346   char        *src, *dst;  in PetscSFLinkGetDataBegin_NVSHMEM()  local
356 …if (direction == PETSCSF_ROOT2LEAF) { /* src is root, dst is leaf; we will move data from src to d…  in PetscSFLinkGetDataBegin_NVSHMEM()
358 …src       = link->rootbuf[PETSCSF_REMOTE][PETSC_MEMTYPE_DEVICE]; /* root buf is the send buf; it i…  in PetscSFLinkGetDataBegin_NVSHMEM()
374   } else { /* src is leaf, dst is root; we will move data from src to dst */  in PetscSFLinkGetDataBegin_NVSHMEM()
376     src       = link->leafbuf[PETSCSF_REMOTE][PETSC_MEMTYPE_DEVICE]; /* leaf buf is the send buf */  in PetscSFLinkGetDataBegin_NVSHMEM()
394   /* After Pack operation -- src tells dst ranks that they are allowed to get data */  in PetscSFLinkGetDataBegin_NVSHMEM()
400   /* dst waits for signals (permissions) from src ranks to start getting data */  in PetscSFLinkGetDataBegin_NVSHMEM()
406 …/* dst gets data from src ranks using non-blocking nvshmem_gets, which are finished in PetscSFLink…  in PetscSFLinkGetDataBegin_NVSHMEM()
408   /* Count number of locally accessible src ranks, which should be a small number */  in PetscSFLinkGetDataBegin_NVSHMEM()
410     if (nvshmem_ptr(src, srcranks_h[i])) nLocallyAccessible++;  in PetscSFLinkGetDataBegin_NVSHMEM()
415 …ssible<<<nsrcranks, 1, 0, link->remoteCommStream>>>(nsrcranks, srcranks_d, src, srcdisp_d, dst, ds…  in PetscSFLinkGetDataBegin_NVSHMEM()
423       if (nvshmem_ptr(src, pe)) {  in PetscSFLinkGetDataBegin_NVSHMEM()
425 …nvshmemx_getmem_nbi_on_stream(dst + (dstdisp_h[i] - dstdisp_h[0]) * link->unitbytes, src + srcdisp…  in PetscSFLinkGetDataBegin_NVSHMEM()
488 …dstranks, PetscMPIInt *dstranks, char *dst, PetscInt *dstdisp, const char *src, PetscInt *srcdisp,…  in WaitAndPutDataToRemotelyAccessible()  argument
497 …nvshmem_putmem_nbi(dst + dstdisp[bid] * unitbytes, src + (srcdisp[bid] - srcdisp[0]) * unitbytes, …  in WaitAndPutDataToRemotelyAccessible()
519   char          *src, *dst;  in PetscSFLinkPutDataBegin_NVSHMEM()  local
530 …src       = link->rootbuf[PETSCSF_REMOTE][PETSC_MEMTYPE_DEVICE]; /* Both src & dst must be symmetr…  in PetscSFLinkPutDataBegin_NVSHMEM()
543     src       = link->leafbuf[PETSCSF_REMOTE][PETSC_MEMTYPE_DEVICE];  in PetscSFLinkPutDataBegin_NVSHMEM()
565 …ks, 1, 0, link->remoteCommStream>>>(ndstranks, dstranks_d, dst, dstdisp_d, src, srcdisp_d, srcsig,…  in PetscSFLinkPutDataBegin_NVSHMEM()
577 …nvshmemx_putmem_nbi_on_stream(dst + dstdisp_h[i] * link->unitbytes, src + (srcdisp_h[i] - srcdisp_…  in PetscSFLinkPutDataBegin_NVSHMEM()
598   /* 2. Wait for signals from src ranks (if any) */  in PutDataEnd()
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()