Lines Matching refs:remoteCommStream

223     PetscCallCUDA(cudaStreamWaitEvent(link->remoteCommStream, link->dataReady, 0));  in PetscSFLinkBuildDependenceBegin()
238     PetscCallCUDA(cudaEventRecord(link->endRemoteComm, link->remoteCommStream));  in PetscSFLinkBuildDependenceEnd()
320 …NvshmemWaitSignals<<<1, 1, 0, link->remoteCommStream>>>(n, sig, 0, 1); /* wait the signals to be 0…  in PetscSFLinkWaitSignalsOfCompletionOfGettingData_NVSHMEM()
396 …NvshmemSendSignals<<<(ndstranks + 255) / 256, 256, 0, link->remoteCommStream>>>(ndstranks, dstsig,…  in PetscSFLinkGetDataBegin_NVSHMEM()
402 …NvshmemWaitSignals<<<1, 1, 0, link->remoteCommStream>>>(nsrcranks, dstsig, 1, 0); /* wait the sign…  in PetscSFLinkGetDataBegin_NVSHMEM()
415 …GetDataFromRemotelyAccessible<<<nsrcranks, 1, 0, link->remoteCommStream>>>(nsrcranks, srcranks_d, …  in PetscSFLinkGetDataBegin_NVSHMEM()
425 …h[0]) * link->unitbytes, src + srcdisp_h[i] * link->unitbytes, nelems, pe, link->remoteCommStream);  in PetscSFLinkGetDataBegin_NVSHMEM()
457 …nvshmemx_quiet_on_stream(link->remoteCommStream); /* Finish the nonblocking get, so that we can un…  in PetscSFLinkGetDataEnd_NVSHMEM()
459 …NvshmemSendSignals<<<(nsrcranks + 511) / 512, 512, 0, link->remoteCommStream>>>(nsrcranks, srcsig,…  in PetscSFLinkGetDataEnd_NVSHMEM()
565 …WaitAndPutDataToRemotelyAccessible<<<ndstranks, 1, 0, link->remoteCommStream>>>(ndstranks, dstrank…  in PetscSFLinkPutDataBegin_NVSHMEM()
571 …WaitSignalsFromLocallyAccessible<<<1, 1, 0, link->remoteCommStream>>>(ndstranks, dstranks_d, srcsi…  in PetscSFLinkPutDataBegin_NVSHMEM()
577 …tbytes, src + (srcdisp_h[i] - srcdisp_h[0]) * link->unitbytes, nelems, pe, link->remoteCommStream);  in PetscSFLinkPutDataBegin_NVSHMEM()
582 …if (nLocallyAccessible) nvshmemx_quiet_on_stream(link->remoteCommStream); /* Calling nvshmem_fence…  in PetscSFLinkPutDataBegin_NVSHMEM()
632 …PutDataEnd<<<1, 1, 0, link->remoteCommStream>>>(nsrcranks, ndstranks, dstranks, dstsig, dstsigdisp…  in PetscSFLinkPutDataEnd_NVSHMEM()
661 …NvshmemSendSignals<<<(nsrcranks + 255) / 256, 256, 0, link->remoteCommStream>>>(nsrcranks, srcsig,…  in PetscSFLinkSendSignalsToAllowPuttingData_NVSHMEM()
675   PetscCallCUDA(cudaStreamDestroy(link->remoteCommStream));  in PetscSFLinkDestroy_NVSHMEM()
762 …PetscCallCUDA(cudaStreamCreateWithPriority(&link->remoteCommStream, cudaStreamNonBlocking, greates…  in PetscSFLinkCreate_NVSHMEM()