.. _memory_model: ============ Memory Model ============ .. _memory_model_diagram: .. image:: img/memory_model-1PPD.png :width: 1000 :align: center An Intel® SHMEM program consists of data objects that are private to each PE and data objects that are remotely accessible by all PEs. Private data objects are stored in the local host memory of a given PE or in the device memory of the device local to the PE, and private objects cannot be accessed by other PEs via ``ishmem`` routines. Private data objects on the host follow the memory model of C/C++ and private data objects on the device follow the SYCL device memory model (see `SYCL device memory model `_). Remotely accessible objects, however, can be accessed by remote PEs using ``ishmem`` routines within SYCL device kernels. Remotely accessible data objects are called `Symmetric Data Objects`. Each symmetric data object has a corresponding object with the same name, type, and size on all PEs where that object is accessible via the Intel® SHMEM API. In Intel® SHMEM, global and static C++ variables are remotely accessible `only` by using ``ishmem`` routines that are callable from the **host** (this kind of data object must not be defined in a dynamic shared object). On the other hand, C++ data allocated by the memory management routines (Section :ref:`Memory Management`) are remotely accessible by using ``ishmem`` APIs that are callable from `either` the **host** or the **device**. .. TODO: .. \footnote{For efficiency reasons, the same offset (from an .. arbitrary memory address) for symmetric data objects might be used on all .. PEs. .. Further discussion about symmetric heap .. layout and implementation efficiency can be found in Section ishmem_free}. .. (For the definition of what is accessible, see the .. descriptions for ishmem_pe_accessible and ishmem_addr_accessible .. in Sections ishmem_pe_accessible and .. ishmem_addr_accessible.) The ``ishmem`` dynamic memory allocation routines (e.g., ``ishmem_malloc``) allow collective allocation of `Symmetric Data Objects` on a special memory region called the `Symmetric Heap`. The Symmetric Heap is created during the execution of a program at a memory location determined by the implementation. The Symmetric Heap may reside in different memory regions on different PEs. :ref:`The memory model diagram` shows an example of the Intel® SHMEM memory layout, illustrating the location of remotely accessible symmetric objects and private data objects. .. _pointers_to_symmetric_objects: ----------------------------- Pointers to Symmetric Objects ----------------------------- Symmetric data objects are referenced in ``ishmem`` routines through the local pointer to the desired remotely accessible object. The address contained in this pointer is referred to as a `symmetric address`. Every symmetric address is also a `local address` that is valid for direct memory access; however, not all local addresses are symmetric. Manipulation of symmetric addresses passed to ``ishmem`` routines - including pointer arithmetic, array indexing, and access of structure or union members - are permitted as long as the resulting local pointer remains within the same symmetric allocation or object. Symmetric addresses are only valid at the PE where they were generated; using a symmetric address generated by a different PE for direct memory access or as an argument to a ``ishmem`` routine results in undefined behavior. Symmetric addresses provided to typed ``ishmem`` interfaces must be naturally aligned based on their type and any requirements of the underlying architecture. .. TODO: .. Symmetric addresses provided to fixed-size Intel® SHMEM interfaces .. (e.g., ``ishmem_put32``) must also be aligned to the given size. .. Symmetric objects provided to fixed-size Intel® SHMEM interfaces .. must have storage size equal to the bit-width of the given .. operation\footnote{The bit-width of a byte is implementation-defined in .. \Cstd. .. The \CONST{CHAR\_BIT} constant in \HEADER{limits.h} can be used to portably .. calculate the bit-width of a \Cstd object.}. .. Because \CorCpp{} structures may contain implementation-defined padding, the .. fixed-size interfaces should not be used with \CorCpp{} structures. .. The `mem` interfaces (e.g., ``ishmem_putmem``) have no alignment .. requirements. The DPC++/SYCL programming model supports Unified Shared Memory between host and GPU. In that model, memory can be `host`, `device`, or `shared`. In Intel® SHMEM, the symmetric heap is in `device` memory and is not accessible from the host by loads and stores, but SYCL routines such as ``sycl::queue::memcpy()`` may be used. SYCL `host` Memory may be allocated by ``sycl::malloc_host()`` and is accessible using loads and stores in both host and device kernel code. SYCL kernels generally do not have access to host global and static objects. Intel® SHMEM augments the DPC++/SYCL memory model to some degree, in that host-initiated operations can reference objects in the symmetric heap on the same or other PEs, even though device memory is not accessible by loads and stores. Host-initiated operations can also refer to global and static objects on the same or other PEs. Device-initiated operations may only reference objects in the symmetric heap. The ``ishmem_ptr`` routine allows the programmer to obtain a local pointer to a remotely accessible object at a specified PE if load-store access is possible. For Intel® :math:`\text{X}^e` Link fabric connected GPUs, ``ishmem_ptr`` may return a valid address for direct access to symmetric heap objects that are accessible via the :math:`\text{X}^e` Link fabric. ``ishmem_ptr`` called on the host will generally not return a valid address for a symmetric heap object even on the same PE. A pointer returned by ``ishmem_ptr`` is valid for direct memory access; however, providing this address as an argument of an ``ishmem`` routine that requires a symmetric address results in undefined behavior. .. _amo_guarantees: -------------------- Atomicity Guarantees -------------------- Intel® SHMEM contains a number of routines that perform atomic operations on symmetric data objects, which are defined in the :ref:`Atomic Memory Operations` section. The atomic routines guarantee that concurrent accesses by any of these routines to the same location, using the same datatype (specified in the :ref:`Standard AMO Types` and :ref:`Extended AMO Types` tables) will be exclusive. Exclusivity is also guaranteed when the target PE performs a wait or test operation on the same location and with the same datatype as one or more atomic operations. .. FIXME: in the same atomicity domain will be exclusive. .. An Intel® SHMEM `atomicity domain` is a set of communication .. contexts whose associated teams (see Section teams) are .. all split by (possibly recursive) calls to a .. ishmem_team_split_* routine from a common predefined team. .. Intel® SHMEM defines two such predefined teams, .. ``ISHMEM_TEAM_WORLD`` and ``ISHMEM_TEAM_SHARED`` (see the `Library .. Handles` section). .. \footnote{ .. Although all PEs in ISHMEM_TEAM_SHARED are also .. in ISHMEM_TEAM_WORLD, and a PE's number can be .. translated from its ISHMEM_TEAM_SHARED to .. ISHMEM_TEAM_WORLD, the .. ISHMEM_TEAM_SHARED team is defined as not having .. been created by a call to a ishmem_team_split_* routine on .. ISHMEM_TEAM_WORLD. .. Therefore, the two teams are distinct predefined teams forming .. separate atomicity domains. .. } Intel® SHMEM atomic operations do not guarantee exclusivity in the following scenarios, all of which result in undefined behavior. #. When concurrent accesses to the same location are performed using ``ishmem`` atomic routines using different datatypes. #. When atomic and non-atomic ``ishmem`` routines are used to access the same location concurrently. #. When ``ishmem`` atomic routines and non-``ishmem`` operations (e.g., load, store, or other atomic operations) are used to access the same location concurrently.