.. _sec:shmem_signal: 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. .. code-block:: c 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); This section specifies the NVSHMEM support for *put-with-signal*, nonblocking *put-with-signal*, and *signal-fetch* routines. The put-with-signal routines provide a method for copying data from a contiguous local data object to a data object on a specified PE and subsequently updating a remote flag to signal completion. The signal-fetch routine provides support for fetching a signal update operation. NVSHMEM *put-with-signal* routines specified in this section have two variants. In one of the variants, the context handle, ``ctx``, is explicitly passed as an argument. In this variant, the operation is performed on the specified context. If the context handle ``ctx`` does not correspond to a valid context, the behavior is undefined. In the other variant, the context handle is not explicitly passed and thus, the operations are performed on the default context. .. _subsec:signal_atomicity: Atomicity Guarantees for Signaling Operations ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ All signaling operations put-with-signal, nonblocking put-with-signal, and signal-fetch are performed on a signal data object, a remotely accessible symmetric object of type ``uint64_t``. A signal operator in the put-with-signal routine is a NVSHMEM library constant that determines the type of update to be performed as a signal on the signal data object. 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; - signal-fetch routine that fetches the signal data object; and - any point-to-point synchronization routine that accesses the signal data object using the same datatype. .. _subsec:signal_operator: Available Signal Operators ~~~~~~~~~~~~~~~~~~~~~~~~~~ With the atomicity guarantees as described in Section :ref:`subsec:signal_atomicity`, the following options can be used as a signal operator. ``NVSHMEM_SIGNAL_SET`` An update to signal data object is an atomic set operation. It writes an unsigned 64-bit value as a signal into the signal data object on a remote ``PE`` as an atomic operation. ``NVSHMEM_SIGNAL_ADD`` An update to signal data object is an atomic add operation. It adds an unsigned 64-bit value as a signal into the signal data object on a remote ``PE`` as an atomic operation. .. _subsec:shmem_put_signal: **NVSHMEM_PUT_SIGNAL** ~~~~~~~~~~~~~~~~~~~~~~ .. c:function:: __device__ void nvshmem\_TYPENAME\_put\_signal(TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe) where *TYPE* is one of the standard RMA types and has a corresponding *TYPENAME* specified by Table :ref:`stdrmatypes`. .. c:function:: __device__ void nvshmem\_putSIZE\_signal(void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe) where *SIZE* is one of ``8, 16, 32, 64, 128``. .. c:function:: __device__ void nvshmem\_putmem\_signal(void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe) *dest [OUT]* Symmetric address of the data object to be updated on the remote PE. The type of ``dest`` should match that implied in the SYNOPSIS section. *source [IN]* Local address of data object containing the data to be copied. The type of ``source`` should match that implied in the SYNOPSIS section. *nelems [IN]* Number of elements in the ``dest`` and ``source`` arrays. For ``nvshmem_putmem_signal`` and ``nvshmem_ctx_putmem_signal``, elements are bytes. *sig_addr [OUT]* Symmetric address of the signal data object to be updated on the remote PE as a signal. *signal [IN]* Unsigned 64-bit value that is used for updating the remote ``sig_addr`` signal data object. *sig_op [IN]* Signal operator that represents the type of update to be performed on the remote ``sig_addr`` signal data object. *pe [IN]* PE number of the remote PE. **Description** The *put-with-signal* routines provide a method for copying data from a contiguous local data object to a data object on a specified PE and subsequently updating a remote flag to signal completion. The routines return after the data has been copied out of the ``source`` array on the local PE. The ``sig_op`` signal operator determines the type of update to be performed on the remote ``sig_addr`` signal data object. The completion of signal update based on the ``sig_op`` signal operator using the ``signal`` flag on the remote PE indicates the delivery of its corresponding ``dest`` data words into the data object on the remote PE. An update to the ``sig_addr`` signal data object through a *put-with-signal* routine completes as if performed atomically as described in Section :ref:`subsec:signal_atomicity`. The various options as described in Section :ref:`subsec:signal_operator` can be used as the ``sig_op`` signal operator. **Returns** None. **Notes** The ``dest`` and ``sig_addr`` data objects must both be remotely accessible. The ``sig_addr`` and ``dest`` could be of different kinds, for example, one could be a global/static *C* variable and the other could be allocated on the symmetric heap. ``sig_addr`` and ``dest`` may not be overlapping in memory. The completion of signal update using the ``signal`` flag on the remote PE indicates only the delivery of its corresponding ``dest`` data words into the data object on the remote PE. Without a memory-ordering operation, there is no implied ordering between the signal update of a *put-with-signal* routine and another data transfer. For example, the completion of the signal update in a sequence consisting of a put routine followed by a *put-with-signal* routine does not imply delivery of the put routine’s data. .. _subsec:shmem_put_signal_nbi: **NVSHMEM_PUT_SIGNAL_NBI** ~~~~~~~~~~~~~~~~~~~~~~~~~~ .. c:function:: __device__ void nvshmem\_TYPENAME\_put\_signal\_nbi(TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe) where *TYPE* is one of the standard RMA types and has a corresponding *TYPENAME* specified by Table :ref:`stdrmatypes`. *dest [OUT]* Symmetric address of the data object to be updated on the remote PE. The type of ``dest`` should match that implied in the SYNOPSIS section. *source [IN]* Local address of data object containing the data to be copied. The type of ``source`` should match that implied in the SYNOPSIS section. *nelems [IN]* Number of elements in the ``dest`` and ``source`` arrays. For ``nvshmem_putmem_signal_nbi`` and ``nvshmem_ctx_putmem_signal_nbi``, elements are bytes. *sig_addr [OUT]* Symmetric address of the signal data object to be updated on the remote PE as a signal. *signal [IN]* Unsigned 64-bit value that is used for updating the remote ``sig_addr`` signal data object. *sig_op [IN]* Signal operator that represents the type of update to be performed on the remote ``sig_addr`` signal data object. *pe [IN]* PE number of the remote PE. **Description** The nonblocking *put-with-signal* routines provide a method for copying data from a contiguous local data object to a data object on a specified PE and subsequently updating a remote flag to signal completion. The routines return after initiating the operation. The operation is considered complete after a subsequent call to ``nvshmem_quiet``. At the completion of ``nvshmem_quiet``, the data has been copied out of the ``source`` array on the local PE and delivered into the ``dest`` array on the destination PE. The delivery of ``signal`` flag on the remote PE indicates only the delivery of its corresponding ``dest`` data words into the data object on the remote PE. Furthermore, two successive nonblocking *put-with-signal* routines, or a nonblocking *put-with-signal* routine with another data transfer may deliver data out of order unless a call to ``nvshmem_fence`` is introduced between the two calls. The ``sig_op`` signal operator determines the type of update to be performed on the remote ``sig_addr`` signal data object. An update to the ``sig_addr`` signal data object through a nonblocking *put-with-signal* routine completes as if performed atomically as described in Section :ref:`subsec:signal_atomicity`. The various options as described in Section :ref:`subsec:signal_operator` can be used as the ``sig_op`` signal operator. **Returns** None. **Notes** The ``dest`` and ``sig_addr`` data objects must both be remotely accessible. The ``sig_addr`` and ``dest`` could be of different kinds, for example, one could be a global/static *C* variable and the other could be allocated on the symmetric heap. ``sig_addr`` and ``dest`` may not be overlapping in memory. .. _subsec:shmem_signal_fetch: **NVSHMEM_SIGNAL_FETCH** ~~~~~~~~~~~~~~~~~~~~~~~~ .. c:function:: __device__ uint64_t nvshmem\_signal\_fetch(const uint64_t *sig_addr) *sig_addr [IN]* Local address of the remotely accessible signal variable. **Description** ``nvshmem_signal_fetch`` performs a fetch operation and returns the contents of the ``sig_addr`` signal data object. Access to ``sig_addr`` signal object at the calling PE is expected to satisfy the atomicity guarantees as described in Section :ref:`subsec:signal_atomicity`. **Returns** Returns the contents of the signal data object, ``sig_addr``, at the calling PE. The following datatypes are supported by the NVSHMEM signal operation. .. _signaltypes: .. table:: Signal Types and Names ================== ========== *TYPE* *TYPENAME* ================== ========== short short int int long long long long longlong unsigned short ushort unsigned int uint unsigned long ulong unsigned long long ulonglong int32_t int32 int64_t int64 uint32_t uint32 uint64_t uint64 size_t size ptrdiff_t ptrdiff ================== ========== .. include:: ../../api/signal.in