Lines Matching refs:link
214 static PetscErrorCode PetscSFLinkBuildDependenceBegin(PetscSF sf, PetscSFLink link, PetscSFDirectio… in PetscSFLinkBuildDependenceBegin() argument
222 PetscCallCUDA(cudaEventRecord(link->dataReady, link->stream)); in PetscSFLinkBuildDependenceBegin()
223 PetscCallCUDA(cudaStreamWaitEvent(link->remoteCommStream, link->dataReady, 0)); in PetscSFLinkBuildDependenceBegin()
229 static PetscErrorCode PetscSFLinkBuildDependenceEnd(PetscSF sf, PetscSFLink link, PetscSFDirection … in PetscSFLinkBuildDependenceEnd() argument
238 PetscCallCUDA(cudaEventRecord(link->endRemoteComm, link->remoteCommStream)); in PetscSFLinkBuildDependenceEnd()
239 PetscCallCUDA(cudaStreamWaitEvent(link->stream, link->endRemoteComm, 0)); in PetscSFLinkBuildDependenceEnd()
304 …FLinkWaitSignalsOfCompletionOfGettingData_NVSHMEM(PetscSF sf, PetscSFLink link, PetscSFDirection d… in PetscSFLinkWaitSignalsOfCompletionOfGettingData_NVSHMEM() argument
312 sig = link->rootSendSig; /* leaf ranks set my rootSendsig */ in PetscSFLinkWaitSignalsOfCompletionOfGettingData_NVSHMEM()
315 sig = link->leafSendSig; in PetscSFLinkWaitSignalsOfCompletionOfGettingData_NVSHMEM()
320 …NvshmemWaitSignals<<<1, 1, 0, link->remoteCommStream>>>(n, sig, 0, 1); /* wait the signals to be 0… in PetscSFLinkWaitSignalsOfCompletionOfGettingData_NVSHMEM()
339 static PetscErrorCode PetscSFLinkGetDataBegin_NVSHMEM(PetscSF sf, PetscSFLink link, PetscSFDirectio… in PetscSFLinkGetDataBegin_NVSHMEM() argument
355 PetscCall(PetscSFLinkBuildDependenceBegin(sf, link, direction)); in PetscSFLinkGetDataBegin_NVSHMEM()
358 …src = link->rootbuf[PETSCSF_REMOTE][PETSC_MEMTYPE_DEVICE]; /* root buf is the send buf; it i… in PetscSFLinkGetDataBegin_NVSHMEM()
366 …dst = link->leafbuf[PETSCSF_REMOTE][PETSC_MEMTYPE_DEVICE]; /* recv buf is the local leaf buf… in PetscSFLinkGetDataBegin_NVSHMEM()
372 dstsig = link->leafRecvSig; in PetscSFLinkGetDataBegin_NVSHMEM()
376 src = link->leafbuf[PETSCSF_REMOTE][PETSC_MEMTYPE_DEVICE]; /* leaf buf is the send buf */ in PetscSFLinkGetDataBegin_NVSHMEM()
384 …dst = link->rootbuf[PETSCSF_REMOTE][PETSC_MEMTYPE_DEVICE]; /* the local root buf is the recv… in PetscSFLinkGetDataBegin_NVSHMEM()
390 dstsig = link->rootRecvSig; in PetscSFLinkGetDataBegin_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 …telyAccessible<<<nsrcranks, 1, 0, link->remoteCommStream>>>(nsrcranks, srcranks_d, src, srcdisp_d,… in PetscSFLinkGetDataBegin_NVSHMEM()
424 size_t nelems = (dstdisp_h[i + 1] - dstdisp_h[i]) * link->unitbytes; in PetscSFLinkGetDataBegin_NVSHMEM()
425 …am(dst + (dstdisp_h[i] - dstdisp_h[0]) * link->unitbytes, src + srcdisp_h[i] * link->unitbytes, ne… in PetscSFLinkGetDataBegin_NVSHMEM()
435 static PetscErrorCode PetscSFLinkGetDataEnd_NVSHMEM(PetscSF sf, PetscSFLink link, PetscSFDirection … in PetscSFLinkGetDataEnd_NVSHMEM() argument
446 srcsig = link->rootSendSig; /* I want to set their root signal */ in PetscSFLinkGetDataEnd_NVSHMEM()
451 srcsig = link->leafSendSig; in PetscSFLinkGetDataEnd_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()
462 PetscCall(PetscSFLinkBuildDependenceEnd(sf, link, direction)); in PetscSFLinkGetDataEnd_NVSHMEM()
514 static PetscErrorCode PetscSFLinkPutDataBegin_NVSHMEM(PetscSF sf, PetscSFLink link, PetscSFDirectio… in PetscSFLinkPutDataBegin_NVSHMEM() argument
527 PetscCall(PetscSFLinkBuildDependenceBegin(sf, link, direction)); in PetscSFLinkPutDataBegin_NVSHMEM()
530 …src = link->rootbuf[PETSCSF_REMOTE][PETSC_MEMTYPE_DEVICE]; /* Both src & dst must be symmetr… in PetscSFLinkPutDataBegin_NVSHMEM()
531 dst = link->leafbuf[PETSCSF_REMOTE][PETSC_MEMTYPE_DEVICE]; in PetscSFLinkPutDataBegin_NVSHMEM()
535 srcsig = link->rootSendSig; in PetscSFLinkPutDataBegin_NVSHMEM()
543 src = link->leafbuf[PETSCSF_REMOTE][PETSC_MEMTYPE_DEVICE]; in PetscSFLinkPutDataBegin_NVSHMEM()
544 dst = link->rootbuf[PETSCSF_REMOTE][PETSC_MEMTYPE_DEVICE]; in PetscSFLinkPutDataBegin_NVSHMEM()
548 srcsig = link->leafSendSig; in PetscSFLinkPutDataBegin_NVSHMEM()
565 …Accessible<<<ndstranks, 1, 0, link->remoteCommStream>>>(ndstranks, dstranks_d, dst, dstdisp_d, src… in PetscSFLinkPutDataBegin_NVSHMEM()
571 …WaitSignalsFromLocallyAccessible<<<1, 1, 0, link->remoteCommStream>>>(ndstranks, dstranks_d, srcsi… in PetscSFLinkPutDataBegin_NVSHMEM()
575 size_t nelems = (srcdisp_h[i + 1] - srcdisp_h[i]) * link->unitbytes; in PetscSFLinkPutDataBegin_NVSHMEM()
577 …_on_stream(dst + dstdisp_h[i] * link->unitbytes, src + (srcdisp_h[i] - srcdisp_h[0]) * link->unitb… in PetscSFLinkPutDataBegin_NVSHMEM()
582 …if (nLocallyAccessible) nvshmemx_quiet_on_stream(link->remoteCommStream); /* Calling nvshmem_fence… in PetscSFLinkPutDataBegin_NVSHMEM()
606 static PetscErrorCode PetscSFLinkPutDataEnd_NVSHMEM(PetscSF sf, PetscSFLink link, PetscSFDirection … in PetscSFLinkPutDataEnd_NVSHMEM() argument
620 dstsig = link->leafRecvSig; /* I will set my leaf ranks's RecvSig */ in PetscSFLinkPutDataEnd_NVSHMEM()
627 dstsig = link->rootRecvSig; in PetscSFLinkPutDataEnd_NVSHMEM()
632 …PutDataEnd<<<1, 1, 0, link->remoteCommStream>>>(nsrcranks, ndstranks, dstranks, dstsig, dstsigdisp… in PetscSFLinkPutDataEnd_NVSHMEM()
635 PetscCall(PetscSFLinkBuildDependenceEnd(sf, link, direction)); in PetscSFLinkPutDataEnd_NVSHMEM()
640 static PetscErrorCode PetscSFLinkSendSignalsToAllowPuttingData_NVSHMEM(PetscSF sf, PetscSFLink link… in PetscSFLinkSendSignalsToAllowPuttingData_NVSHMEM() argument
650 srcsig = link->rootSendSig; /* I want to set their send signals */ in PetscSFLinkSendSignalsToAllowPuttingData_NVSHMEM()
655 srcsig = link->leafSendSig; in PetscSFLinkSendSignalsToAllowPuttingData_NVSHMEM()
661 …NvshmemSendSignals<<<(nsrcranks + 255) / 256, 256, 0, link->remoteCommStream>>>(nsrcranks, srcsig,… in PetscSFLinkSendSignalsToAllowPuttingData_NVSHMEM()
668 static PetscErrorCode PetscSFLinkDestroy_NVSHMEM(PetscSF sf, PetscSFLink link) in PetscSFLinkDestroy_NVSHMEM() argument
673 PetscCallCUDA(cudaEventDestroy(link->dataReady)); in PetscSFLinkDestroy_NVSHMEM()
674 PetscCallCUDA(cudaEventDestroy(link->endRemoteComm)); in PetscSFLinkDestroy_NVSHMEM()
675 PetscCallCUDA(cudaStreamDestroy(link->remoteCommStream)); in PetscSFLinkDestroy_NVSHMEM()
678 PetscCall(PetscNvshmemFree(link->leafbuf_alloc[PETSCSF_REMOTE][PETSC_MEMTYPE_DEVICE])); in PetscSFLinkDestroy_NVSHMEM()
679 PetscCall(PetscNvshmemFree(link->leafSendSig)); in PetscSFLinkDestroy_NVSHMEM()
680 PetscCall(PetscNvshmemFree(link->leafRecvSig)); in PetscSFLinkDestroy_NVSHMEM()
681 PetscCall(PetscNvshmemFree(link->rootbuf_alloc[PETSCSF_REMOTE][PETSC_MEMTYPE_DEVICE])); in PetscSFLinkDestroy_NVSHMEM()
682 PetscCall(PetscNvshmemFree(link->rootSendSig)); in PetscSFLinkDestroy_NVSHMEM()
683 PetscCall(PetscNvshmemFree(link->rootRecvSig)); in PetscSFLinkDestroy_NVSHMEM()
691 PetscSFLink *p, link; in PetscSFLinkCreate_NVSHMEM() local
721 for (p = &bas->avail; (link = *p); p = &link->next) { in PetscSFLinkCreate_NVSHMEM()
722 if (link->use_nvshmem) { in PetscSFLinkCreate_NVSHMEM()
723 PetscCall(MPIPetsc_Type_compare(unit, link->unit, &match)); in PetscSFLinkCreate_NVSHMEM()
725 *p = link->next; /* Remove from available list */ in PetscSFLinkCreate_NVSHMEM()
730 PetscCall(PetscNew(&link)); in PetscSFLinkCreate_NVSHMEM()
731 …PetscCall(PetscSFLinkSetUp_Host(sf, link, unit)); /* Comp… in PetscSFLinkCreate_NVSHMEM()
732 …if (sf->backend == PETSCSF_BACKEND_CUDA) PetscCall(PetscSFLinkSetUp_CUDA(sf, link, unit)); /* Setu… in PetscSFLinkCreate_NVSHMEM()
734 …else if (sf->backend == PETSCSF_BACKEND_KOKKOS) PetscCall(PetscSFLinkSetUp_Kokkos(sf, link, unit)); in PetscSFLinkCreate_NVSHMEM()
737 …link->rootdirect[PETSCSF_LOCAL] = PETSC_TRUE; /* For the local part we directly use root/leafdata … in PetscSFLinkCreate_NVSHMEM()
738 link->leafdirect[PETSCSF_LOCAL] = PETSC_TRUE; in PetscSFLinkCreate_NVSHMEM()
741 …if (!link->rootSendSig) PetscCall(PetscNvshmemCalloc(bas->nRemoteLeafRanksMax * sizeof(uint64_t), … in PetscSFLinkCreate_NVSHMEM()
742 …if (!link->rootRecvSig) PetscCall(PetscNvshmemCalloc(bas->nRemoteLeafRanksMax * sizeof(uint64_t), … in PetscSFLinkCreate_NVSHMEM()
743 …if (!link->leafSendSig) PetscCall(PetscNvshmemCalloc(sf->nRemoteRootRanksMax * sizeof(uint64_t), (… in PetscSFLinkCreate_NVSHMEM()
744 …if (!link->leafRecvSig) PetscCall(PetscNvshmemCalloc(sf->nRemoteRootRanksMax * sizeof(uint64_t), (… in PetscSFLinkCreate_NVSHMEM()
746 link->use_nvshmem = PETSC_TRUE; in PetscSFLinkCreate_NVSHMEM()
747 link->rootmtype = PETSC_MEMTYPE_DEVICE; /* Only need 0/1-based mtype from now on */ in PetscSFLinkCreate_NVSHMEM()
748 link->leafmtype = PETSC_MEMTYPE_DEVICE; in PetscSFLinkCreate_NVSHMEM()
750 link->Destroy = PetscSFLinkDestroy_NVSHMEM; in PetscSFLinkCreate_NVSHMEM()
752 link->PrePack = PetscSFLinkWaitSignalsOfCompletionOfGettingData_NVSHMEM; in PetscSFLinkCreate_NVSHMEM()
753 link->StartCommunication = PetscSFLinkGetDataBegin_NVSHMEM; in PetscSFLinkCreate_NVSHMEM()
754 link->FinishCommunication = PetscSFLinkGetDataEnd_NVSHMEM; in PetscSFLinkCreate_NVSHMEM()
756 link->StartCommunication = PetscSFLinkPutDataBegin_NVSHMEM; in PetscSFLinkCreate_NVSHMEM()
757 link->FinishCommunication = PetscSFLinkPutDataEnd_NVSHMEM; in PetscSFLinkCreate_NVSHMEM()
758 link->PostUnpack = PetscSFLinkSendSignalsToAllowPuttingData_NVSHMEM; in PetscSFLinkCreate_NVSHMEM()
762 …PetscCallCUDA(cudaStreamCreateWithPriority(&link->remoteCommStream, cudaStreamNonBlocking, greates… in PetscSFLinkCreate_NVSHMEM()
764 PetscCallCUDA(cudaEventCreateWithFlags(&link->dataReady, cudaEventDisableTiming)); in PetscSFLinkCreate_NVSHMEM()
765 PetscCallCUDA(cudaEventCreateWithFlags(&link->endRemoteComm, cudaEventDisableTiming)); in PetscSFLinkCreate_NVSHMEM()
769 …link->rootbuf[PETSCSF_REMOTE][PETSC_MEMTYPE_DEVICE] = (char *)rootdata + bas->rootstart[PETSCSF_RE… in PetscSFLinkCreate_NVSHMEM()
771 …if (!link->rootbuf_alloc[PETSCSF_REMOTE][PETSC_MEMTYPE_DEVICE]) PetscCall(PetscNvshmemMalloc(bas->… in PetscSFLinkCreate_NVSHMEM()
772 …link->rootbuf[PETSCSF_REMOTE][PETSC_MEMTYPE_DEVICE] = link->rootbuf_alloc[PETSCSF_REMOTE][PETSC_ME… in PetscSFLinkCreate_NVSHMEM()
776 …link->leafbuf[PETSCSF_REMOTE][PETSC_MEMTYPE_DEVICE] = (char *)leafdata + sf->leafstart[PETSCSF_REM… in PetscSFLinkCreate_NVSHMEM()
778 …if (!link->leafbuf_alloc[PETSCSF_REMOTE][PETSC_MEMTYPE_DEVICE]) PetscCall(PetscNvshmemMalloc(sf->l… in PetscSFLinkCreate_NVSHMEM()
779 …link->leafbuf[PETSCSF_REMOTE][PETSC_MEMTYPE_DEVICE] = link->leafbuf_alloc[PETSCSF_REMOTE][PETSC_ME… in PetscSFLinkCreate_NVSHMEM()
782 link->rootdirect[PETSCSF_REMOTE] = rootdirect[PETSCSF_REMOTE]; in PetscSFLinkCreate_NVSHMEM()
783 link->leafdirect[PETSCSF_REMOTE] = leafdirect[PETSCSF_REMOTE]; in PetscSFLinkCreate_NVSHMEM()
784 …link->rootdata = rootdata; /* root/leafdata are keys to look up links in PetscSF… in PetscSFLinkCreate_NVSHMEM()
785 link->leafdata = leafdata; in PetscSFLinkCreate_NVSHMEM()
786 link->next = bas->inuse; in PetscSFLinkCreate_NVSHMEM()
787 bas->inuse = link; in PetscSFLinkCreate_NVSHMEM()
788 *mylink = link; in PetscSFLinkCreate_NVSHMEM()