.. _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);

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.

.. _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
   ================== ==========

.. include:: ../../api/signal.in