Lines Matching refs:sig
253 __global__ static void NvshmemSendSignals(PetscInt n, uint64_t *sig, PetscInt *sigdisp, PetscMPIInt… in NvshmemSendSignals() argument
258 if (i < n) nvshmemx_uint64_signal(sig + sigdisp[i], newval, ranks[i]); in NvshmemSendSignals()
269 __global__ static void NvshmemWaitSignals(PetscInt n, uint64_t *sig, uint64_t expval, uint64_t newv… in NvshmemWaitSignals() argument
275 nvshmem_signal_wait_until(sig+i,NVSHMEM_CMP_EQ,expval); in NvshmemWaitSignals()
276 sig[i] = newval; in NvshmemWaitSignals()
279 nvshmem_uint64_wait_until_all(sig, n, NULL /*no mask*/, NVSHMEM_CMP_EQ, expval); in NvshmemWaitSignals()
280 for (int i = 0; i < n; i++) sig[i] = newval; in NvshmemWaitSignals()
307 uint64_t *sig; in PetscSFLinkWaitSignalsOfCompletionOfGettingData_NVSHMEM() local
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()