Searched refs:remoteCommStream (Results 1 – 2 of 2) sorted by relevance
| /petsc/src/vec/is/sf/impls/basic/nvshmem/ |
| H A D | sfnvshmem.cu | 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()
[all …]
|
| /petsc/src/vec/is/sf/impls/basic/ |
| H A D | sfpack.h | 202 cupmStream_t remoteCommStream; /* Streams for remote (i.e., inter-rank) communication */ member
|