327 __global__ static void GetDataFromRemotelyAccessible(PetscInt nsrcranks, PetscMPIInt *srcranks, con…  in GetDataFromRemotelyAccessible()  argument
344   PetscInt nsrcranks, ndstranks, nLocallyAccessible = 0;  in PetscSFLinkGetDataBegin_NVSHMEM()  local
357     nsrcranks = sf->nRemoteRootRanks;  in PetscSFLinkGetDataBegin_NVSHMEM()
375     nsrcranks = bas->nRemoteLeafRanks;  in PetscSFLinkGetDataBegin_NVSHMEM()
401   if (nsrcranks) {  in PetscSFLinkGetDataBegin_NVSHMEM()
402 …NvshmemWaitSignals<<<1, 1, 0, link->remoteCommStream>>>(nsrcranks, dstsig, 1, 0); /* wait the sign…  in PetscSFLinkGetDataBegin_NVSHMEM()
409   for (int i = 0; i < nsrcranks; i++) {  in PetscSFLinkGetDataBegin_NVSHMEM()
414   if (nLocallyAccessible < nsrcranks) {  in PetscSFLinkGetDataBegin_NVSHMEM()
415 …GetDataFromRemotelyAccessible<<<nsrcranks, 1, 0, link->remoteCommStream>>>(nsrcranks, srcranks_d, …  in PetscSFLinkGetDataBegin_NVSHMEM()
421     for (int i = 0; i < nsrcranks; i++) {  in PetscSFLinkGetDataBegin_NVSHMEM()
440   PetscInt       nsrcranks, *srcsigdisp;  in PetscSFLinkGetDataEnd_NVSHMEM()  local
445     nsrcranks  = sf->nRemoteRootRanks;  in PetscSFLinkGetDataEnd_NVSHMEM()
450     nsrcranks  = bas->nRemoteLeafRanks;  in PetscSFLinkGetDataEnd_NVSHMEM()
456   if (nsrcranks) {  in PetscSFLinkGetDataEnd_NVSHMEM()
459 …NvshmemSendSignals<<<(nsrcranks + 511) / 512, 512, 0, link->remoteCommStream>>>(nsrcranks, srcsig,…  in PetscSFLinkGetDataEnd_NVSHMEM()
587 __global__ static void PutDataEnd(PetscInt nsrcranks, PetscInt ndstranks, PetscMPIInt *dstranks, ui…  in PutDataEnd()  argument
599   if (nsrcranks) {  in PutDataEnd()
600 …nvshmem_uint64_wait_until_all(dstsig, nsrcranks, NULL /*no mask*/, NVSHMEM_CMP_EQ, 1); /* wait sig…  in PutDataEnd()
601     for (int i = 0; i < nsrcranks; i++) dstsig[i] = 0;  in PutDataEnd()
612   PetscInt       nsrcranks, ndstranks, *dstsigdisp;  in PetscSFLinkPutDataEnd_NVSHMEM()  local
616     nsrcranks = sf->nRemoteRootRanks;  in PetscSFLinkPutDataEnd_NVSHMEM()
623     nsrcranks = bas->nRemoteLeafRanks;  in PetscSFLinkPutDataEnd_NVSHMEM()
631   if (nsrcranks || ndstranks) {  in PetscSFLinkPutDataEnd_NVSHMEM()
632 …PutDataEnd<<<1, 1, 0, link->remoteCommStream>>>(nsrcranks, ndstranks, dstranks, dstsig, dstsigdisp…  in PetscSFLinkPutDataEnd_NVSHMEM()
644   PetscInt       nsrcranks, *srcsigdisp_d;  in PetscSFLinkSendSignalsToAllowPuttingData_NVSHMEM()  local
649     nsrcranks    = sf->nRemoteRootRanks;  in PetscSFLinkSendSignalsToAllowPuttingData_NVSHMEM()
654     nsrcranks    = bas->nRemoteLeafRanks;  in PetscSFLinkSendSignalsToAllowPuttingData_NVSHMEM()
660   if (nsrcranks) {  in PetscSFLinkSendSignalsToAllowPuttingData_NVSHMEM()
661 …NvshmemSendSignals<<<(nsrcranks + 255) / 256, 256, 0, link->remoteCommStream>>>(nsrcranks, srcsig,…  in PetscSFLinkSendSignalsToAllowPuttingData_NVSHMEM()