Signaling Operations¶
NVSHMEM provides signaling operations that can be used to update a remote flag variable. When used in conjunction with wait/test routines at the remote PE, these routines can provide efficient point-to-point synchronization. The following example shows signal operations used to implement neighbor communication in a ring.
nvshmem_putmem(dest, src, size, (pe+1) % npes);
nvshmem_quiet();
nvshmemx_int_signal(flag, 1, (pe+1) % npes);
nvshmem_int_wait_until(flag, NVSHMEM_CMP_EQ, 1);
Atomicity Guarantees for Signaling Operations [subsec:signal_atomicity]
All signaling operations complete as if performed atomically with respect to the following:
- other signal operations that update the signal data object using the same datatype;
- any point-to-point synchronization routine that accesses the signal data object using the same datatype.
The following datatypes are supported by the NVSHMEM signal operation.
| TYPE | TYPENAME |
|---|---|
| short | short |
| int | int |
| long | long |
| long long | longlong |
| unsigned short | ushort |
| unsigned int | uint |
| unsigned long | ulong |
| unsigned long long | ulonglong |
NVSHMEMX_SIGNAL¶
-
__device__ inline void
nvshmemx_TYPENAME_signal(TYPE *dest, const TYPE value, int pe)¶
where TYPE is one of the standard RMA types and has a corresponding TYPENAME specified by Table Signal Types and Names.
- dest [OUT]
- Symmetric address of the signal word to be updated.
- value [IN]
- The value to be placed in dest.
- pe [IN]
- PE number of the remote PE.
Description
The nvshmemx_signal operation atomically sets dest to value on the specified PE. This operation can be used together with wait and test routines for efficient point-to-point synchronization.
Returns
None.