.. _sec:amo:

Atomic Memory Operations
------------------------

An AMO is a one-sided communication mechanism that combines memory read,
update, or write operations with atomicity guarantees described in
Section  :ref:`subsec:amo_guarantees`. Similar
to the RMA routines, described in Section :ref:`sec:rma`, the
AMOs are performed only on symmetric objects. NVSHMEM defines two types
of AMO routines:

-  The *fetching* routines return the original value of, and optionally
   update, the remote data object in a single atomic operation. The
   routines return after the data has been fetched from the target PE
   and delivered to the calling PE. The data type of the returned value
   is the same as the type of the remote data object.

-  The *non-fetching* routines update the remote data object in a single
   atomic operation. A call to a non-fetching atomic routine issues the
   atomic operation and may return before the operation executes on the
   target PE. The ``nvshmem_quiet``, ``nvshmem_barrier``, or
   ``nvshmem_barrier_all`` routines can be used to force completion for
   these non-fetching atomic routines.

   The non-fetching routines include:

NVSHMEM provides AMO interfaces with the following types:

.. _stdamotypes:

.. table:: Standard AMO Types and Names

   ========= ==========
   *TYPE*    *TYPENAME*
   ========= ==========
   int       int
   long      long
   long long longlong
   size_t    size
   ptrdiff_t ptrdiff
   ========= ==========

.. _extamotypes:

.. table:: Extended AMO Types and Names

   ========= ==========
   *TYPE*    *TYPENAME*
   ========= ==========
   float     float
   double    double
   int       int
   long      long
   long long longlong
   size_t    size
   ptrdiff_t ptrdiff
   ========= ==========

.. _bitamotypes:

.. table:: Bitwise AMO Types and Names

   ================== ==========
   *TYPE*             *TYPENAME*
   ================== ==========
   unsigned int       uint
   unsigned long      ulong
   unsigned long long ulonglong
   int32_t            int32
   int64_t            int64
   uint32_t           uint32
   uint64_t           uint64
   ================== ==========

.. _subsec:shmem_atomic_fetch:

**NVSHMEM_ATOMIC_FETCH**
^^^^^^^^^^^^^^^^^^^^^^^^



.. c:function:: TYPE nvshmem\_TYPENAME\_atomic\_fetch(const TYPE *source, int pe)



.. c:function:: __device__ TYPE nvshmem\_TYPENAME\_atomic\_fetch(const TYPE *source, int pe)

where *TYPE* is one of the extended AMO types and has a corresponding
*TYPENAME* specified by Table :ref:`extamotypes`.



.. c:function:: TYPE nvshmem\_fetch(const TYPE *source, int pe)



.. c:function:: __device__ TYPE nvshmem\_fetch(const TYPE *source, int pe)

where *TYPE* is one of {``float``, ``double``, ``int``, ``long``,
``long long``}.



.. c:function:: TYPE nvshmem\_TYPENAME\_fetch(const TYPE *source, int pe)



.. c:function:: __device__ TYPE nvshmem\_TYPENAME\_fetch(const TYPE *source, int pe)

where *TYPE* is one of {``float``, ``double``, ``int``, ``long``,
``long long``} and has a corresponding *TYPENAME* specified by
Table :ref:`extamotypes`.

*source [IN]*
   Symmetric address of the source data object. The type of ``source``
   should match that implied in the SYNOPSIS section.

*pe [IN]*
   An integer that indicates the PE number from which ``source`` is to
   be fetched.

**Description**

``nvshmem_atomic_fetch`` performs an atomic fetch operation. It returns
the contents of the ``source`` as an atomic operation.

**Returns**

The contents at the ``source`` address on the remote PE. The data type
of the return value is the same as the type of the remote data object.

.. _subsec:shmem_atomic_set:

**NVSHMEM_ATOMIC_SET**
^^^^^^^^^^^^^^^^^^^^^^



.. c:function:: void nvshmem\_TYPENAME\_atomic\_set(TYPE *dest, TYPE value, int pe)



.. c:function:: __device__ void nvshmem\_TYPENAME\_atomic\_set(TYPE *dest, TYPE value, int pe)

where *TYPE* is one of the extended AMO types and has a corresponding
*TYPENAME* specified by Table :ref:`extamotypes`.



.. c:function:: void nvshmem\_set(TYPE *dest, TYPE value, int pe)



.. c:function:: __device__ void nvshmem\_set(TYPE *dest, TYPE value, int pe)

where *TYPE* is one of {``float``, ``double``, ``int``, ``long``,
``long long``}.



.. c:function:: void nvshmem\_TYPENAME\_set(TYPE *dest, TYPE value, int pe)



.. c:function:: __device__ void nvshmem\_TYPENAME\_set(TYPE *dest, TYPE value, int pe)

where *TYPE* is one of {``float``, ``double``, ``int``, ``long``,
``long long``} and has a corresponding *TYPENAME* specified by
Table :ref:`extamotypes`.

*dest [OUT]*
   Symmetric address of the destination data object. The type of
   ``dest`` should match that implied in the SYNOPSIS section.

*value [IN]*
   The operand to the atomic set operation. The type of ``value`` should
   match that implied in the SYNOPSIS section.

*pe [IN]*
   An integer that indicates the PE number on which ``dest`` is to be
   updated.

**Description**

``nvshmem_atomic_set`` performs an atomic set operation. It writes the
``value`` into ``dest`` on ``pe`` as an atomic operation.

**Returns**

None.

.. _subsec:shmem_atomic_compare_swap:

**NVSHMEM_ATOMIC_COMPARE_SWAP**
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^



.. c:function:: TYPE nvshmem\_TYPENAME\_atomic\_compare\_swap(TYPE *dest, TYPE cond, TYPE value, int pe)



.. c:function:: __device__ TYPE nvshmem\_TYPENAME\_atomic\_compare\_swap(TYPE *dest, TYPE cond, TYPE value, int pe)

where *TYPE* is one of the standard AMO types and has a corresponding
*TYPENAME* specified by Table :ref:`stdamotypes`.



.. c:function:: TYPE nvshmem\_cswap(TYPE *dest, TYPE cond, TYPE value, int pe)



.. c:function:: __device__ TYPE nvshmem\_cswap(TYPE *dest, TYPE cond, TYPE value, int pe)

where *TYPE* is one of {``int``, ``long``, ``long long``}.



.. c:function:: TYPE nvshmem\_TYPENAME\_cswap(TYPE *dest, TYPE cond, TYPE value, int pe)



.. c:function:: __device__ TYPE nvshmem\_TYPENAME\_cswap(TYPE *dest, TYPE cond, TYPE value, int pe)

where *TYPE* is one of {``int``, ``long``, ``long long``} and has a
corresponding *TYPENAME* specified by
Table :ref:`stdamotypes`.

*dest [OUT]*
   Symmetric address of the destination data object. The type of
   ``dest`` should match that implied in the SYNOPSIS section.

*cond [IN]*
   ``cond`` is compared to the remote ``dest`` value. If ``cond`` and
   the remote ``dest`` are equal, then ``value`` is swapped into the
   remote ``dest``; otherwise, the remote ``dest`` is unchanged. In
   either case, the old value of the remote ``dest`` is returned as the
   routine return value. ``cond`` must be of the same data type as
   ``dest``.

*value [IN]*
   The value to be atomically written to the remote PE. The type of
   ``value`` should match that implied in the SYNOPSIS section.

*pe [IN]*
   An integer that indicates the PE number upon which ``dest`` is to be
   updated.

**Description**

The conditional swap routines conditionally update a ``dest`` data
object on the specified PE and return the prior contents of the data
object in one atomic operation.

**Returns**

The contents that had been in the ``dest`` data object on the remote PE
prior to the conditional swap. Data type is the same as the ``dest``
data type.

.. _subsec:shmem_atomic_swap:

**NVSHMEM_ATOMIC_SWAP**
^^^^^^^^^^^^^^^^^^^^^^^



.. c:function:: TYPE nvshmem\_TYPENAME\_atomic\_swap(TYPE *dest, TYPE value, int pe)



.. c:function:: __device__ TYPE nvshmem\_TYPENAME\_atomic\_swap(TYPE *dest, TYPE value, int pe)

where *TYPE* is one of the extended AMO types and has a corresponding
*TYPENAME* specified by Table :ref:`extamotypes`.



.. c:function:: TYPE nvshmem\_swap(TYPE *dest, TYPE value, int pe)



.. c:function:: __device__ TYPE nvshmem\_swap(TYPE *dest, TYPE value, int pe)

where *TYPE* is one of {``float``, ``double``, ``int``, ``long``,
``long long``}.



.. c:function:: TYPE nvshmem\_TYPENAME\_swap(TYPE *dest, TYPE value, int pe)



.. c:function:: __device__ TYPE nvshmem\_TYPENAME\_swap(TYPE *dest, TYPE value, int pe)

where *TYPE* is one of {``float``, ``double``, ``int``, ``long``,
``long long``} and has a corresponding *TYPENAME* specified by
Table :ref:`extamotypes`.

*dest [OUT]*
   Symmetric address of the destination data object. The type of
   ``dest`` should match that implied in the SYNOPSIS section.

*value [IN]*
   The value to be atomically written to the remote PE. The type of
   ``value`` should match that implied in the SYNOPSIS section.

*pe [IN]*
   An integer that indicates the PE number on which ``dest`` is to be
   updated.

**Description**

``nvshmem_atomic_swap`` performs an atomic swap operation. It writes
``value`` into ``dest`` on PE and returns the previous contents of
``dest`` as an atomic operation.

**Returns**

The content that had been at the ``dest`` address on the remote PE prior
to the swap is returned.

.. _subsec:shmem_atomic_fetch_inc:

**NVSHMEM_ATOMIC_FETCH_INC**
^^^^^^^^^^^^^^^^^^^^^^^^^^^^



.. c:function:: TYPE nvshmem\_TYPENAME\_atomic\_fetch\_inc(TYPE *dest, int pe)



.. c:function:: __device__ TYPE nvshmem\_TYPENAME\_atomic\_fetch\_inc(TYPE *dest, int pe)

where *TYPE* is one of the standard AMO types and has a corresponding
*TYPENAME* specified by Table :ref:`stdamotypes`.



.. c:function:: TYPE nvshmem\_finc(TYPE *dest, int pe)



.. c:function:: __device__ TYPE nvshmem\_finc(TYPE *dest, int pe)

where *TYPE* is one of {``int``, ``long``, ``long long``}.



.. c:function:: TYPE nvshmem\_TYPENAME\_finc(TYPE *dest, int pe)



.. c:function:: __device__ TYPE nvshmem\_TYPENAME\_finc(TYPE *dest, int pe)

where *TYPE* is one of {``int``, ``long``, ``long long``} and has a
corresponding *TYPENAME* specified by
Table :ref:`stdamotypes`.

*dest [OUT]*
   Symmetric address of the destination data object. The type of
   ``dest`` should match that implied in the SYNOPSIS section.

*pe [IN]*
   An integer that indicates the PE number on which ``dest`` is to be
   updated.

**Description**

These routines perform a fetch-and-increment operation. The ``dest`` on
PE ``pe`` is increased by one and the routine returns the previous
contents of ``dest`` as an atomic operation.

**Returns**

The contents that had been at the ``dest`` address on the remote PE
prior to the increment. The data type of the return value is the same as
the ``dest``.

.. _subsec:shmem_atomic_inc:

**NVSHMEM_ATOMIC_INC**
^^^^^^^^^^^^^^^^^^^^^^



.. c:function:: void nvshmem\_TYPENAME\_atomic\_inc(TYPE *dest, int pe)



.. c:function:: __device__ void nvshmem\_TYPENAME\_atomic\_inc(TYPE *dest, int pe)

where *TYPE* is one of the standard AMO types and has a corresponding
*TYPENAME* specified by Table :ref:`stdamotypes`.



.. c:function:: void nvshmem\_inc(TYPE *dest, int pe)



.. c:function:: __device__ void nvshmem\_inc(TYPE *dest, int pe)

where *TYPE* is one of {``int``, ``long``, ``long long``}.



.. c:function:: void nvshmem\_TYPENAME\_inc(TYPE *dest, int pe)



.. c:function:: __device__ void nvshmem\_TYPENAME\_inc(TYPE *dest, int pe)

where *TYPE* is one of {``int``, ``long``, ``long long``} and has a
corresponding *TYPENAME* specified by
Table :ref:`stdamotypes`.

*dest [OUT]*
   Symmetric address of the destination data object. The type of
   ``dest`` should match that implied in the SYNOPSIS section.

*pe [IN]*
   An integer that indicates the PE number on which ``dest`` is to be
   updated.

**Description**

These routines perform an atomic increment operation on the ``dest``
data object on PE.

**Returns**

None.

.. _subsec:shmem_atomic_fetch_add:

**NVSHMEM_ATOMIC_FETCH_ADD**
^^^^^^^^^^^^^^^^^^^^^^^^^^^^



.. c:function:: TYPE nvshmem\_TYPENAME\_atomic\_fetch\_add(TYPE *dest, TYPE value, int pe)



.. c:function:: __device__ TYPE nvshmem\_TYPENAME\_atomic\_fetch\_add(TYPE *dest, TYPE value, int pe)

where *TYPE* is one of the standard AMO types and has a corresponding
*TYPENAME* specified by Table :ref:`stdamotypes`.



.. c:function:: TYPE nvshmem\_fadd(TYPE *dest, TYPE value, int pe)



.. c:function:: __device__ TYPE nvshmem\_fadd(TYPE *dest, TYPE value, int pe)

where *TYPE* is one of {``int``, ``long``, ``long long``}.



.. c:function:: TYPE nvshmem\_TYPENAME\_fadd(TYPE *dest, TYPE value, int pe)



.. c:function:: __device__ TYPE nvshmem\_TYPENAME\_fadd(TYPE *dest, TYPE value, int pe)

where *TYPE* is one of {``int``, ``long``, ``long long``} and has a
corresponding *TYPENAME* specified by
Table :ref:`stdamotypes`.

*dest [OUT]*
   Symmetric address of the destination data object. The type of
   ``dest`` should match that implied in the SYNOPSIS section.

*value [IN]*
   The operand to the atomic fetch-and-add operation. The type of
   ``value`` should match that implied in the SYNOPSIS section.

*pe [IN]*
   An integer that indicates the PE number on which ``dest`` is to be
   updated.

**Description**

``nvshmem_atomic_fetch_add`` routines perform an atomic fetch-and-add
operation. An atomic fetch-and-add operation fetches the old ``dest``
and adds ``value`` to ``dest`` without the possibility of another atomic
operation on the ``dest`` between the time of the fetch and the update.
These routines add ``value`` to ``dest`` on ``pe`` and return the
previous contents of ``dest`` as an atomic operation.

**Returns**

The contents that had been at the ``dest`` address on the remote PE
prior to the atomic addition operation. The data type of the return
value is the same as the ``dest``.

.. _subsec:shmem_atomic_add:

**NVSHMEM_ATOMIC_ADD**
^^^^^^^^^^^^^^^^^^^^^^



.. c:function:: void nvshmem\_TYPENAME\_atomic\_add(TYPE *dest, TYPE value, int pe)



.. c:function:: __device__ void nvshmem\_TYPENAME\_atomic\_add(TYPE *dest, TYPE value, int pe)

where *TYPE* is one of the standard AMO types and has a corresponding
*TYPENAME* specified by Table :ref:`stdamotypes`.



.. c:function:: void nvshmem\_add(TYPE *dest, TYPE value, int pe)



.. c:function:: __device__ void nvshmem\_add(TYPE *dest, TYPE value, int pe)

where *TYPE* is one of {``int``, ``long``, ``long long``}.



.. c:function:: void nvshmem\_TYPENAME\_add(TYPE *dest, TYPE value, int pe)



.. c:function:: __device__ void nvshmem\_TYPENAME\_add(TYPE *dest, TYPE value, int pe)

where *TYPE* is one of {``int``, ``long``, ``long long``} and has a
corresponding *TYPENAME* specified by
Table :ref:`stdamotypes`.

*dest [OUT]*
   Symmetric address of the destination data object. The type of
   ``dest`` should match that implied in the SYNOPSIS section.

*value [IN]*
   The operand to the atomic add operation. The type of ``value`` should
   match that implied in the SYNOPSIS section.

*pe [IN]*
   An integer that indicates the PE number upon which ``dest`` is to be
   updated.

**Description**

The ``nvshmem_atomic_add`` routine performs an atomic add operation. It
adds ``value`` to ``dest`` on PE ``pe`` and atomically updates the
``dest`` without returning the value.

**Returns**

None.

.. _subsec:shmem_atomic_fetch_and:

**NVSHMEM_ATOMIC_FETCH_AND**
^^^^^^^^^^^^^^^^^^^^^^^^^^^^



.. c:function:: TYPE nvshmem\_TYPENAME\_atomic\_fetch\_and(TYPE *dest, TYPE value, int pe)



.. c:function:: __device__ TYPE nvshmem\_TYPENAME\_atomic\_fetch\_and(TYPE *dest, TYPE value, int pe)

where *TYPE* is one of the bitwise AMO types and has a corresponding
*TYPENAME* specified by Table :ref:`bitamotypes`.

*dest [OUT]*
   Symmetric address of the destination data object. The type of
   ``dest`` should match that implied in the SYNOPSIS section.

*value [IN]*
   The operand to the bitwise AND operation. The type of ``value``
   should match that implied in the SYNOPSIS section.

*pe [IN]*
   An integer value for the PE on which ``dest`` is to be updated.

**Description**

``nvshmem_atomic_fetch_and`` atomically performs a fetching bitwise AND
on the remotely accessible data object pointed to by ``dest`` at PE
``pe`` with the operand ``value``.

**Returns**

The value pointed to by ``dest`` on PE ``pe`` immediately before the
operation is performed.

.. _subsec:shmem_atomic_and:

**NVSHMEM_ATOMIC_AND**
^^^^^^^^^^^^^^^^^^^^^^



.. c:function:: void nvshmem\_TYPENAME\_atomic\_and(TYPE *dest, TYPE value, int pe)



.. c:function:: __device__ void nvshmem\_TYPENAME\_atomic\_and(TYPE *dest, TYPE value, int pe)

where *TYPE* is one of the bitwise AMO types and has a corresponding
*TYPENAME* specified by Table :ref:`bitamotypes`.

*dest [OUT]*
   Symmetric address of the destination data object. The type of
   ``dest`` should match that implied in the SYNOPSIS section.

*value [IN]*
   The operand to the bitwise AND operation. The type of ``value``
   should match that implied in the SYNOPSIS section.

*pe [IN]*
   An integer value for the PE on which ``dest`` is to be updated.

**Description**

``nvshmem_atomic_and`` atomically performs a non-fetching bitwise AND on
the remotely accessible data object pointed to by ``dest`` at PE ``pe``
with the operand ``value``.

**Returns**

None.

.. _subsec:shmem_atomic_fetch_or:

**NVSHMEM_ATOMIC_FETCH_OR**
^^^^^^^^^^^^^^^^^^^^^^^^^^^



.. c:function:: TYPE nvshmem\_TYPENAME\_atomic\_fetch\_or(TYPE *dest, TYPE value, int pe)



.. c:function:: __device__ TYPE nvshmem\_TYPENAME\_atomic\_fetch\_or(TYPE *dest, TYPE value, int pe)

where *TYPE* is one of the bitwise AMO types and has a corresponding
*TYPENAME* specified by Table :ref:`bitamotypes`.

*dest [OUT]*
   Symmetric address of the destination data object. The type of
   ``dest`` should match that implied in the SYNOPSIS section.

*value [IN]*
   The operand to the bitwise OR operation. The type of ``value`` should
   match that implied in the SYNOPSIS section.

*pe [IN]*
   An integer value for the PE on which ``dest`` is to be updated.

**Description**

``nvshmem_atomic_fetch_or`` atomically performs a fetching bitwise OR on
the remotely accessible data object pointed to by ``dest`` at PE ``pe``
with the operand ``value``.

**Returns**

The value pointed to by ``dest`` on PE ``pe`` immediately before the
operation is performed.

.. _subsec:shmem_atomic_or:

**NVSHMEM_ATOMIC_OR**
^^^^^^^^^^^^^^^^^^^^^



.. c:function:: void nvshmem\_TYPENAME\_atomic\_or(TYPE *dest, TYPE value, int pe)



.. c:function:: __device__ void nvshmem\_TYPENAME\_atomic\_or(TYPE *dest, TYPE value, int pe)

where *TYPE* is one of the bitwise AMO types and has a corresponding
*TYPENAME* specified by Table :ref:`bitamotypes`.

*dest [OUT]*
   Symmetric address of the destination data object. The type of
   ``dest`` should match that implied in the SYNOPSIS section.

*value [IN]*
   The operand to the bitwise OR operation. The type of ``value`` should
   match that implied in the SYNOPSIS section.

*pe [IN]*
   An integer value for the PE on which ``dest`` is to be updated.

**Description**

``nvshmem_atomic_or`` atomically performs a non-fetching bitwise OR on
the remotely accessible data object pointed to by ``dest`` at PE ``pe``
with the operand ``value``.

**Returns**

None.

.. _subsec:shmem_atomic_fetch_xor:

**NVSHMEM_ATOMIC_FETCH_XOR**
^^^^^^^^^^^^^^^^^^^^^^^^^^^^



.. c:function:: TYPE nvshmem\_TYPENAME\_atomic\_fetch\_xor(TYPE *dest, TYPE value, int pe)



.. c:function:: __device__ TYPE nvshmem\_TYPENAME\_atomic\_fetch\_xor(TYPE *dest, TYPE value, int pe)

where *TYPE* is one of the bitwise AMO types and has a corresponding
*TYPENAME* specified by Table :ref:`bitamotypes`.

*dest [OUT]*
   Symmetric address of the destination data object. The type of
   ``dest`` should match that implied in the SYNOPSIS section.

*value [IN]*
   The operand to the bitwise XOR operation. The type of ``value``
   should match that implied in the SYNOPSIS section.

*pe [IN]*
   An integer value for the PE on which ``dest`` is to be updated.

**Description**

``nvshmem_atomic_fetch_xor`` atomically performs a fetching bitwise XOR
on the remotely accessible data object pointed to by ``dest`` at PE
``pe`` with the operand ``value``.

**Returns**

The value pointed to by ``dest`` on PE ``pe`` immediately before the
operation is performed.

.. _subsec:shmem_atomic_xor:

**NVSHMEM_ATOMIC_XOR**
^^^^^^^^^^^^^^^^^^^^^^



.. c:function:: void nvshmem\_TYPENAME\_atomic\_xor(TYPE *dest, TYPE value, int pe)



.. c:function:: __device__ void nvshmem\_TYPENAME\_atomic\_xor(TYPE *dest, TYPE value, int pe)

where *TYPE* is one of the bitwise AMO types and has a corresponding
*TYPENAME* specified by Table :ref:`bitamotypes`.

*dest [OUT]*
   Symmetric address of the destination data object. The type of
   ``dest`` should match that implied in the SYNOPSIS section.

*value [IN]*
   The operand to the bitwise XOR operation. The type of ``value``
   should match that implied in the SYNOPSIS section.

*pe [IN]*
   An integer value for the PE on which ``dest`` is to be updated.

**Description**

``nvshmem_atomic_xor`` atomically performs a non-fetching bitwise XOR on
the remotely accessible data object pointed to by ``dest`` at PE ``pe``
with the operand ``value``.

**Returns**

None.